Skip to content

SYCL: add oneMKL GEMM flash attention for XMX-accelerated prompt proc… - #25025

Merged
ggerganov merged 18 commits into
ggml-org:masterfrom
johnkarlhill:sycl-mkl-flash-attn
Jul 31, 2026
Merged

SYCL: add oneMKL GEMM flash attention for XMX-accelerated prompt proc…#25025
ggerganov merged 18 commits into
ggml-org:masterfrom
johnkarlhill:sycl-mkl-flash-attn

Conversation

@johnkarlhill

@johnkarlhill johnkarlhill commented Jun 26, 2026

Copy link
Copy Markdown
Contributor

Adds a flash attention path that routes Q·K^T and S·V matrix multiplies
through oneMKL GEMM, enabling XMX hardware acceleration on Intel GPUs.

Motivation

The existing SYCL flash attention kernels (VEC, TILE) run entirely in
SYCL subgroup operations. On Intel Arc GPUs with XMX matrix engines
(Battlemage and later), oneMKL GEMM can process the large matmuls in
attention significantly faster — particularly at high context lengths.
where the KV cache is quantized.

When it activates

- KV cache is quantized (q8_0, q4_0, q4_1, q5_0, q5_1, or any K-quant)

  • Fast Attention is enabled (--flash-attn on or -fa)
  • K sequence length ≥ 1024 tokens (covers the full --batch-size)
    - Q sequence length ≥ 128
  • Q sequence length ≥ 32 (routes all multi-token prefill through MKL
    while keeping single-token decode and MTP spec drafts on the existing
    TG-optimized VEC kernel)

~~These thresholds route prompt processing through MKL while leaving
single-token decode to the existing TG-optimized kernels. ~~
The path is never activated for f16/bf16 KV cache — those already perform well with the TILE kernel and graph capture.

All KV cache types benefit from XMX acceleration — F16 (the default),
BF16, F32, and quantized (q8_0, q4_0, q4_1, q5_0, q5_1, K-quant).
The MKL kernel converts non-F16 K/V to F16 via to_fp16_sycl before
GEMM, so no additional conversion code was needed.

Set GGML_SYCL_ENABLE_MKL_FA=0 to force the TILE/VEC path for A/B
testing or power comparison.

Implementation

All logic is in one new file, fattn-mkl.cpp (567 lines). The pipeline:

  1. Dequantize K/V to fp16
  2. For each KV head: pack all GQA query heads into a single fp16 buffer
  3. Chunked KV loop (8192-token chunks):
    • MKL GEMM: KQ = Q_batched × K_chunk^T
    • Online softmax SYCL kernel (row-wise, with running max/sum)
    • MKL GEMM: VKQ_chunk = S × V_chunk
    • Accumulate: VKQ_accum += VKQ_chunk
  4. Normalize each GQA head by KQ_sum and scatter to output

GQA groups sharing a KV head are batched into single GEMM calls —
6 query heads × 1020 tokens = 6120 rows in one MKL call, amortizing
launch overhead.

All benchmarks: Qwen 3.6 27B UD Q5_K_XL, MTP enabled.

Context KV Cache PP t/s TG t/s
32K f16 (default) ~671 ~15
32K bf16 ~668 ~12
32K q8_0 ~671 ~15
110K q8_0 ~335 ~17

For comparison, stock bf16 KV cache + FA off on the same GPU achieves 822 t/s PP at 8K — the MKL path with q8_0 is within 1% while using quantized memory.
The f16/bf16 numbers are with FA-on and the MKL path — matching the
quantized-cache baseline that previously required --cache-type-k q8_0

Testing

  • test-backend-ops: 3605/3605 FLASH_ATTN_EXT tests pass (all
    quant types, head sizes 64–512, causal/non-causal masks, sinks,
    max_bias, GQA ratios, multi-batch)
  • Multi-batch: parallel-2 at 32K context, stable throughput,
    no coherence errors
  • Cache reuse: full-context slot restore with LCP similarity,
    MKL path correctly handles the reprocess delta
  • Build isolation: all code is SYCL-only, gated behind
    BEST_FATTN_KERNEL_MKL enum; other backends and non-quantized
    paths are completely unaffected

Known limitations

  • No graph capture: MKL GEMM's internal queue management is
    incompatible with SYCL command graph replay. The existing
    GGML_SYCL_DISABLE_GRAPH default (1) handles this.
  • No ALiBi: max_bias == 0.0f asserted; models needing ALiBi
    will fall through to the TILE kernel.
  • No sinks tensor: dst->src[4] is not yet supported.

Debug output

Timing instrumentation is gated behind MKL_FA_DEBUG=1. In normal operation the MKL path produces no output.
Timing instrumentation is gated behind GGML_SYCL_MKL_FA_DEBUG=1.
Output fingerprint for correctness verification is available with
GGML_SYCL_MKL_FA_DIAG=1. In normal operation the MKL path produces
no output.

AI disclosure

Claude Code was used for SYCL boilerplate (ND-range kernel launches,
ggml_sycl_pool_alloc patterns) and initial drafting of the chunked KV
loop. All algorithmic decisions — oneMKL GEMM integration, online
softmax with GQA batching, activation thresholds, chunk sizing — were
human-directed. Comprehensive testing (3605 test-backend-ops,
multi-quant and multi-batch coherence validation, performance
benchmarking at contexts up to 110K) was performed manually.


🤖 Generated with Claude Code using DeepSeek-V4-Pro

@johnkarlhill
johnkarlhill requested a review from a team as a code owner June 26, 2026 01:24
@github-actions github-actions Bot added ggml changes relating to the ggml tensor library for machine learning SYCL https://en.wikipedia.org/wiki/SYCL - GPU programming language labels Jun 26, 2026
@ggml-gh-bot

ggml-gh-bot Bot commented Jun 26, 2026

Copy link
Copy Markdown

Hi @johnkarlhill, thanks for your contribution!

Per our contribution guidelines, the automated PR checker found the following issue(s) that need your attention:

  • AI-generated content: This project does not accept PRs, descriptions or commit messages that are fully or predominantly AI-generated. If you have used AI to assist you in writing code, please make sure to disclose that explicitly.

Please note that maintainers reserve the right to make final decisions on PRs. If you believe there is a mistake, please comment below.

@arthw arthw left a comment

Copy link
Copy Markdown
Contributor

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

@johnkarlhill

It's good to see this PR to enable XMX in FA.

Could you share which LLM show good performance increasing by this PR?
I use Qwen3.6 and can't trigger the oneMKL path on FA.

Thank you!

@johnkarlhill

Copy link
Copy Markdown
Contributor Author

Adding before and after... both compiled with arch flags to show side-by-side. Compiling without arch flags will degrade performance from these numbers but should still be better than stock.
PR25025 - Qwen3.6-27B-MTP-UD-Q5_K_XL on B70.txt
b9752 - Qwen3.6-27B-MTP-UD-Q5_K_XL on B70.txt

I'll add more models if needed.

@arthw

arthw commented Jun 26, 2026

Copy link
Copy Markdown
Contributor

@johnkarlhill
1.
Could you provide a smaller LLM case to show the perf increase for this PR?
Including the whole cmd.

For user, how to trigger the new code in usage?

Thank you!

@maxious

maxious commented Jun 26, 2026

Copy link
Copy Markdown
Contributor

Tested on dual Intel Arc Pro B60 (Battlemage, 24GB each), oneAPI 2026.0, MKL 2026.0, targeting bmg-g31 AOT.

Short context (pp512 — MKL path NOT active)

No regression vs master:

Model Size Master pp512 PR pp512
gpt-oss 20B Q8_0 11.3G 854 851

Long context (pp2048 — MKL path active)

Model Size Master pp2048 PR pp2048 Delta
llama-2 7B Q2_K 2.6G 950 1,102 +16%
Llama-3 8B Q4_0 4.3G 888 968 +9%
gpt-oss 20B Q8_0 11.3G 503 504 0%
Qwen3.6-35B-A3B MoE Q3_K 12.8G 575 575 0%

Decode speeds (tg128) unchanged across all models (±2%).

Full commands

# Master
./build-master/bin/llama-bench -m model.gguf -p 512,2048 -n 128 -ngl -1 -fa 1

# PR #25025
./build-pr/bin/llama-bench -m model.gguf -p 512,2048 -n 128 -ngl -1 -fa 1

@johnkarlhill

johnkarlhill commented Jun 26, 2026

Copy link
Copy Markdown
Contributor Author

Updated arch table:

GPU -DGGML_SYCL_DEVICE_ARCH
Arc A770 / A750 acm_g10
Arc A580 acm_g12
Arc A380 / A310 acm_g11
Arc B580 / B570 / Pro B70 bmg_g21
Flex / Data Center Max pvc
Integrated (Meteor Lake) mtl_u
Integrated (Lunar Lake) lnl_m

All of these use underscores (_), never hyphens (-).

I incorrectly listed Arc A770 / A750 as acm_g12.

"For user, how to trigger the new code in usage?"
Use "--batch-size N" where N is a value >= 1024.

@jlionhan

jlionhan commented Jun 26, 2026

Copy link
Copy Markdown

I hope this helps.

255H, Arc 140T, 32GB RAM, llama-cli

Build options:

cmake --fresh -B build -DCMAKE_C_COMPILER=icx -DCMAKE_CXX_COMPILER=icpx -DCMAKE_BUILD_TYPE=Release -DGGML_SYCL=1 -DBUILD_SHARED_LIBS=0 -DGGML_SYCL_F16=1 

after PR:

model size params backend ngl threads type_k type_v fa test t/s
gemma4 26B.A4B Q5_K - Medium 17.80 GiB 25.23 B SYCL 99 6 q8_0 q8_0 1 pp512 176.84 ± 3.13
gemma4 26B.A4B Q5_K - Medium 17.80 GiB 25.23 B SYCL 99 6 q8_0 q8_0 1 pp1024 202.22 ± 2.48
gemma4 26B.A4B Q5_K - Medium 17.80 GiB 25.23 B SYCL 99 6 q8_0 q8_0 1 pp2048 212.19 ± 1.95
gemma4 26B.A4B Q5_K - Medium 17.80 GiB 25.23 B SYCL 99 6 q8_0 q8_0 1 tg128 12.13 ± 0.22

before PR:

model size params backend ngl threads type_k type_v fa test t/s
gemma4 26B.A4B Q5_K - Medium 17.80 GiB 25.23 B SYCL 99 6 q8_0 q8_0 1 pp512 178.96 ± 6.26
gemma4 26B.A4B Q5_K - Medium 17.80 GiB 25.23 B SYCL 99 6 q8_0 q8_0 1 pp1024 165.24 ± 3.03
gemma4 26B.A4B Q5_K - Medium 17.80 GiB 25.23 B SYCL 99 6 q8_0 q8_0 1 pp2048 139.39 ± 1.68
gemma4 26B.A4B Q5_K - Medium 17.80 GiB 25.23 B SYCL 99 6 q8_0 q8_0 1 tg128 12.02 ± 0.38

build: f728ada (9793)

Thank you.

However, I am observing intermittent behavior where the most recently entered prompt is not being processed, and the model instead generates a response to the previous prompt. I believe further testing is needed to confirm whether this issue is reproducible and to identify the underlying cause. This behavior may be unrelated to this PR.

@johnkarlhill

johnkarlhill commented Jun 26, 2026

Copy link
Copy Markdown
Contributor Author

However, I am observing intermittent behavior where the most recently entered prompt is not being processed, and the model instead generates a response to the previous prompt. I believe further testing is needed to confirm whether this issue is reproducible and to identify the underlying cause. This behavior may be unrelated to this PR.

I can reproduce the behavior on Gemma4 models and working on a fix. This behavior does not exist on Qwen models. Multiple folks have tested with a few different Qwen models and can't reproduce this. It seems specific to Gemma4.

And a huge THANK YOU for testing this. It is very much appreciated!!!

@arthw

arthw commented Jun 27, 2026

Copy link
Copy Markdown
Contributor

@johnkarlhill
There are several code to call wait().
Are they necessary to get the correct result?
Reduce or remove them will be quicker.

Thank you!

@johnkarlhill

johnkarlhill commented Jun 28, 2026

Copy link
Copy Markdown
Contributor Author

Bug fix

The MKL normalize kernel was writing output using a dense head-major layout (head * n_queries * DV), but llama.cpp's flash attention output uses an interleaved layout (query * n_heads + head per row, matching TILE's flash_attn_combine_results). Head 0 row 0 happened to alias at offset 0 in both layouts, so the first layer's first 64 output floats matched TILE. Everything else landed at wrong addresses. One-line fix in mkl_fa_normalize_head.

Tested models (all pass multi-turn coherence)

  • Gemma-4-26B-A4B-it (Q5_K_M, gqa=2)
  • Gemma-4-31B-it-qat (Q4_K_XL, dense)
  • Qwen3.6-27B (Q5_K_XL, gqa=6)
  • Qwen3.6-35B-A3B (Q4_K_XL, gqa=2)

Performance (Intel Arc Pro B70, Battlemage BMG-G21, 32K context, q8_0 KV cache)

Model MKL PP (t/s) TILE PP (t/s) Speedup
Gemma-4-26B 1473 746 1.97×
Qwen3.6-27B 606 330 1.84×

Token generation unaffected (±1 t/s, within noise) — MKL only activates for prompt processing (n_kv ≥ 1024 with quantized KV).

Updated PR25025 - Qwen3.6-27B-MTP-UD-Q5_K_XL on B70.txt
Updated PR25025 - Gemma-4-26B-A4B-it-UD-Q5_K_M on B70.txt

How to test

cmake --preset x64-windows-sycl-release -DGGML_SYCL_F16=ON -DGGML_SYCL_DEVICE_ARCH=bmg_g21
cmake --build build-x64-windows-sycl-release --config Release -j 16

# Run (MKL activates automatically with flash-attn + quantized KV + n_kv ≥ 1024)
llama-server --flash-attn on --cache-type-k q8_0 --cache-type-v q8_0 --batch-size 1024 ...

# Disable for A/B comparison
set MKL_FA_DISABLE=1```

@johnkarlhill
johnkarlhill force-pushed the sycl-mkl-flash-attn branch from a4871d8 to ce37155 Compare June 28, 2026 16:42
@johnkarlhill

Copy link
Copy Markdown
Contributor Author

@johnkarlhill There are several code to call wait(). Are they necessary to get the correct result? Reduce or remove them will be quicker.

Thank you!

Removed 7 redundant stream->wait() calls — 3 no-ops covered by the SYCL in-order queue, 3 in a diagnostic block that was already gated behind MKL_FA_DIAG=1 (removed the whole block), and 1 final drain before pool destructors that's unnecessary with in-order semantics.

The 4 remaining waits are all required: oneMKL gemm() runs on its own internal queue that does not respect the SYCL in-order queue. Without these barriers, the softmax kernel would read stale KQ data and the accumulate kernel would read stale VKQ_chunk data. TILE uses zero explicit waits because everything is pure SYCL — MKL can't avoid these handshake points.

Perf unchanged. 609 (new) vs 606 t/s... within noise.

- Fix mkl_fa_normalize_head: use interleaved dst layout
  ((query * n_q_heads + head) * DV) matching TILE's
  flash_attn_combine_results. Previously used dense head-major
  layout which wrote head outputs to wrong addresses, corrupting
  attention for all models except Qwen3.6-27B (where GQA=6 heads
  were sparse enough to avoid visible overlap).

- Remove 7 redundant stream->wait() calls — SYCL in-order queue
  already serializes pure SYCL kernel dependencies. Retain only
  the 4 MKL GEMM ↔ SYCL handshake barriers (oneMKL GEMM uses its
  own internal queue that does not respect SYCL in-order).

- Remove unused dst_row_stride, diagnostic clutter, and dead
  K/V hex dump (fa_diag block in fattn-mkl.cpp).

- Add MKL_FA_DISABLE=1 env var for A/B testing.
- Add FA-DISP watchdog (MKL_FA_DEBUG=1) and FA-DIAG output
  fingerprint (MKL_FA_DIAG=1) in fattn.cpp.

Tested: Gemma-4-26B, Gemma-4-31B, Qwen3.6-27B, Qwen3.6-35B-A3B
Perf (B70/Battlemage, 32K, q8_0 KV):
  Gemma-4-26B:  1473 t/s MKL vs 746 TILE (1.97x)
  Qwen3.6-27B:   609 t/s MKL vs 330 TILE (1.85x)

Co-Authored-By: Claude Code on DeepSeek-v4-Pro
Comment thread ggml/src/ggml-sycl/fattn.cpp Outdated
static int fa_diag = -1;
static int fa_diag_count = 0;
if (fa_diag < 0) {
const char * e = getenv("MKL_FA_DIAG");

Copy link
Copy Markdown
Contributor

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

Rename:
MKL_FA_DISABLE to GGML_SYCL_ENABLE_MKL_FA.
MKL_FA_DIAG to GGML_SYCL_MKL_FA_DIAG

Explain them in SYCL.md, refer to chapter: # Environment Variable.

Comment thread ggml/src/ggml-sycl/fattn.cpp Outdated
if (kb == BEST_FATTN_KERNEL_MKL) kname = "MKL";
if (kb == BEST_FATTN_KERNEL_TILE) kname = "TILE";
if (kb == BEST_FATTN_KERNEL_VEC) kname = "VEC";
fprintf(stderr, "[FA-DIAG] #%d %s D=%d n_kv=%lld n_q=%lld "

Copy link
Copy Markdown
Contributor

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

fprintf(stderr, ) is replaced by GGML_LOG_INFO()

Comment thread ggml/src/ggml-sycl/fattn.cpp Outdated
// the same D — helps detect cache-truncation issues.
static int nkv_debug = -1;
if (nkv_debug < 0) {
const char * e = getenv("MKL_FA_DEBUG");

Copy link
Copy Markdown
Contributor

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

Rename:
MKL_FA_DEBUG to GGML_SYCL_MKL_FA_DEBUG
Explain it in SYCL.md

Comment thread ggml/src/ggml-sycl/fattn.cpp Outdated
const char * e = getenv("MKL_FA_DISABLE");
mkl_disable = (e && e[0] == '1') ? 1 : 0;
}
if (mkl_disable == 0 && Q->ne[1] >= 128 && K->ne[1] >= 1024

Copy link
Copy Markdown
Contributor

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

    if (mkl_disable == 0 && Q->ne[1] >= 128 && K->ne[1] >= 1024

Add comment to explain how to trigger FA_MKL in usage or llama-cli/server/bench parameters.

…, document in SYCL.md

Completed the following:
- Rename MKL_FA_DISABLE → GGML_SYCL_ENABLE_MKL_FA (inverted: 0 to disable)
- Rename MKL_FA_DEBUG → GGML_SYCL_MKL_FA_DEBUG
- Rename MKL_FA_DIAG → GGML_SYCL_MKL_FA_DIAG
- Replace fprintf(stderr, ...) / fflush(stderr) with GGML_LOG_INFO() macro
- Document all three env vars in docs/backend/SYCL.md under Runtime
- Add comment explaining MKL FA activation trigger (flash-attn + quantized
  KV cache + batch-size >= 1024 + n_kv >= 1024)

Resolves review feedback from arthw.
Again, thank you!!!

Co-Authored-By: Claude Code on DeepSeek-v4-Pro
@johnkarlhill
johnkarlhill force-pushed the sycl-mkl-flash-attn branch from ce37155 to 5a81c11 Compare June 29, 2026 03:10
@github-actions github-actions Bot added the documentation Improvements or additions to documentation label Jun 29, 2026
Comment thread ggml/src/ggml-sycl/fattn-mkl.cpp Outdated
// MIT license
// Copyright (C) 2025 Intel Corporation
// SPDX-License-Identifier: MIT
//

Copy link
Copy Markdown
Contributor

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

llama.cpp follow the unified copyright definition.
So, no need to declare here.

remove:

//
// MIT license
// Copyright (C) 2025 Intel Corporation
// SPDX-License-Identifier: MIT
//

Comment thread docs/backend/SYCL.md Outdated
| GGML_SYCL_USE_LEVEL_ZERO_API | 1 (default) or 0 | Use Level Zero API for device memory allocation instead of SYCL. Reduces system RAM usage on Intel dGPUs by avoiding DMA-buf/TTM host memory staging. Requires GGML_SYCL_SUPPORT_LEVEL_ZERO_API=ON at build time. SYCL backend always runs on Level Zero running time even if it's set as OFF (The SYCL api will be usage for memory allocation).|
| GGML_SYCL_DISABLE_DNN | 0 (default) or 1 | Disable running computations through oneDNN and always use oneMKL. |
| GGML_SYCL_ENABLE_VMM | 0 or 1 (default) | Enable the virtual-memory device pool. |
| GGML_SYCL_ENABLE_MKL_FA | 1 (default) or 0 | Enable oneMKL GEMM flash attention for XMX-accelerated prompt processing with quantized KV cache. Set to 0 to force the TILE kernel path for A/B testing. Activated at runtime when flash-attn is enabled (`-fa` or `--flash-attn on`), KV cache is quantized (`--cache-type-k/q *_0/*_1`), and KV length ≥ 1024. |

Copy link
Copy Markdown
Contributor

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

User can find the parameter: --cache-type-k in llama-cli.

But how to set for KV length ≥ 1024?
Please provide a detailed method in llama-cli/server for common user.

Comment thread ggml/src/ggml-sycl/fattn-mkl.cpp Outdated
Comment on lines +327 to +328
#define MKL_TAKE_TIME(t0) auto t0 = std::chrono::steady_clock::now()
#define MKL_ACCUM(acc, t0) acc += (int64_t)std::chrono::duration_cast \

Copy link
Copy Markdown
Contributor

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

The two macros are used to debug for perf.
They should be disabled as default.

Comment thread ggml/src/ggml-sycl/fattn-mkl.cpp Outdated
Comment on lines +569 to +570
stream->wait();
try { ev.wait_and_throw(); } catch (sycl::exception & e) {

Copy link
Copy Markdown
Contributor

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

stream->wait() & ev.wait_and_throw() are duplicated wait() code.
Maybe impact side effect.
Remove one of them.

Comment thread ggml/src/ggml-sycl/fattn-mkl.cpp Outdated
Comment on lines +603 to +604
stream->wait();
try { ev.wait_and_throw(); } catch (sycl::exception & e) {

Copy link
Copy Markdown
Contributor

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

same comment as above: remove one.

Comment thread ggml/src/ggml-sycl/fattn.cpp Outdated
Comment on lines +135 to +136
const char * e = getenv("GGML_SYCL_ENABLE_MKL_FA");
mkl_disable = (e && e[0] == '0') ? 1 : 0;

Copy link
Copy Markdown
Contributor

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

The code can be replaced by existed function, like
ggml_sycl_get_env("GGML_SYCL_USM_SYSTEM", 0);

Comment thread ggml/src/ggml-sycl/fattn.cpp Outdated
Comment thread ggml/src/ggml-sycl/fattn.cpp Outdated
…ove dup waits, gate perf macros

- Replace raw getenv() with ggml_sycl_get_env() in all 4 env-var checks
  (fattn.cpp: GGML_SYCL_ENABLE_MKL_FA, GGML_SYCL_MKL_FA_DEBUG,
   GGML_SYCL_MKL_FA_DIAG; fattn-mkl.cpp: GGML_SYCL_MKL_FA_DEBUG)
- Remove duplicated stream->wait() before ev.wait_and_throw() in GEMM
  KQ and GEMM VKQ — ev.wait_and_throw() already waits for completion
- Gate MKL_ACCUM macro behind do_print so timing accumulators are
  no-ops in normal operation
- Remove redundant MIT/Intel copyright header from fattn-mkl.cpp
- Remove unused #include <cfloat>
- Expand SYCL.md MKL FA docs with step-by-step activation trigger
  and example llama-cli command

Again, thank you!!!

Co-Authored-By: Claude Code on DeepSeek-v4-Pro
Remove the quantized-only restriction on MKL activation — the MKL
kernel converts any non-F16 K/V to F16 via to_fp16_sycl before GEMM,
so F16 (default), BF16, and F32 caches all benefit from XMX hardware
acceleration.  The type restriction was an unnecessary gate.

Before (F16/BF16 default cache + FA on at 32K prefill): ~356 t/s (TILE path)
After:  ~670 t/s (MKL path, matching quantized-cache baseline)

Minimal change: two conditions removed, one comment updated in fattn.cpp.
No kernel or conversion code changes — the dequant pipeline already
covers all types.
@johnkarlhill

Copy link
Copy Markdown
Contributor Author

@arthw Updated based on your suggestions. Thank you!!

@arthw arthw left a comment

Copy link
Copy Markdown
Contributor

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

It's good job!

Thank you!

@arthw

arthw commented Jul 13, 2026

Copy link
Copy Markdown
Contributor

@johnkarlhill
Please resolve the conflict!
Thank you!

@johnkarlhill

Copy link
Copy Markdown
Contributor Author

@arthw I think you meant the conflict with the new GGML_SYCL_ENABLE_FUSION env variable, so I added that. Please let me know if there was another conflict I needed to address.

@NeoZhangJianyu

Copy link
Copy Markdown
Contributor

@johnkarlhill
The conflict means the github show conflict in the page.
It disappear now.

Thank you!

@arthw arthw added the merge ready A maintainer can use this label to indicate that they consider the changes final and ready to merge. label Jul 14, 2026
@ggerganov
ggerganov merged commit 9d9a6d2 into ggml-org:master Jul 31, 2026
26 of 28 checks passed
huaxel pushed a commit to huaxel/CachyLLama that referenced this pull request Aug 2, 2026
ggml-org#25025)

* SYCL: add oneMKL GEMM flash attention for XMX-accelerated prompt processing

* fattn-mkl: fix interleaved dst layout in normalize kernel

- Fix mkl_fa_normalize_head: use interleaved dst layout
  ((query * n_q_heads + head) * DV) matching TILE's
  flash_attn_combine_results. Previously used dense head-major
  layout which wrote head outputs to wrong addresses, corrupting
  attention for all models except Qwen3.6-27B (where GQA=6 heads
  were sparse enough to avoid visible overlap).

- Remove 7 redundant stream->wait() calls — SYCL in-order queue
  already serializes pure SYCL kernel dependencies. Retain only
  the 4 MKL GEMM ↔ SYCL handshake barriers (oneMKL GEMM uses its
  own internal queue that does not respect SYCL in-order).

- Remove unused dst_row_stride, diagnostic clutter, and dead
  K/V hex dump (fa_diag block in fattn-mkl.cpp).

- Add MKL_FA_DISABLE=1 env var for A/B testing.
- Add FA-DISP watchdog (MKL_FA_DEBUG=1) and FA-DIAG output
  fingerprint (MKL_FA_DIAG=1) in fattn.cpp.

Tested: Gemma-4-26B, Gemma-4-31B, Qwen3.6-27B, Qwen3.6-35B-A3B
Perf (B70/Battlemage, 32K, q8_0 KV):
  Gemma-4-26B:  1473 t/s MKL vs 746 TILE (1.97x)
  Qwen3.6-27B:   609 t/s MKL vs 330 TILE (1.85x)

Co-Authored-By: Claude Code on DeepSeek-v4-Pro

* Thank you for the review feedback: rename env vars, use GGML_LOG_INFO, document in SYCL.md

Completed the following:
- Rename MKL_FA_DISABLE → GGML_SYCL_ENABLE_MKL_FA (inverted: 0 to disable)
- Rename MKL_FA_DEBUG → GGML_SYCL_MKL_FA_DEBUG
- Rename MKL_FA_DIAG → GGML_SYCL_MKL_FA_DIAG
- Replace fprintf(stderr, ...) / fflush(stderr) with GGML_LOG_INFO() macro
- Document all three env vars in docs/backend/SYCL.md under Runtime
- Add comment explaining MKL FA activation trigger (flash-attn + quantized
  KV cache + batch-size >= 1024 + n_kv >= 1024)

Resolves review feedback from arthw.
Again, thank you!!!

Co-Authored-By: Claude Code on DeepSeek-v4-Pro

* Thank you for the review feedback round 2: use ggml_sycl_get_env, remove dup waits, gate perf macros

- Replace raw getenv() with ggml_sycl_get_env() in all 4 env-var checks
  (fattn.cpp: GGML_SYCL_ENABLE_MKL_FA, GGML_SYCL_MKL_FA_DEBUG,
   GGML_SYCL_MKL_FA_DIAG; fattn-mkl.cpp: GGML_SYCL_MKL_FA_DEBUG)
- Remove duplicated stream->wait() before ev.wait_and_throw() in GEMM
  KQ and GEMM VKQ — ev.wait_and_throw() already waits for completion
- Gate MKL_ACCUM macro behind do_print so timing accumulators are
  no-ops in normal operation
- Remove redundant MIT/Intel copyright header from fattn-mkl.cpp
- Remove unused #include <cfloat>
- Expand SYCL.md MKL FA docs with step-by-step activation trigger
  and example llama-cli command

Again, thank you!!!

Co-Authored-By: Claude Code on DeepSeek-v4-Pro

* fattn-mkl: enable MKL FA for all KV cache types

Remove the quantized-only restriction on MKL activation — the MKL
kernel converts any non-F16 K/V to F16 via to_fp16_sycl before GEMM,
so F16 (default), BF16, and F32 caches all benefit from XMX hardware
acceleration.  The type restriction was an unnecessary gate.

Before (F16/BF16 default cache + FA on at 32K prefill): ~356 t/s (TILE path)
After:  ~670 t/s (MKL path, matching quantized-cache baseline)

Minimal change: two conditions removed, one comment updated in fattn.cpp.
No kernel or conversion code changes — the dequant pipeline already
covers all types.

* fattn-mkl: rename mkl_disable -> mkl_enable for clarity

* fattn-mkl: refine MKL FA dispatch gates

Three changes:
1. Remove quantized-only restriction - MKL FA activates for all
   KV cache types (F16 default, BF16, F32, quantized).  The MKL
   kernel converts non-F16 K/V via to_fp16_sycl before GEMM.
2. Rename mkl_disable -> mkl_enable to match env var
   (GGML_SYCL_ENABLE_MKL_FA).
3. Replace batch-size threshold with Q->ne[1] >= 32 gate.
   Keeps TG (Q=1) and MTP drafts (Q=3-8) on VEC path where
   fused kernel beats MKL launch overhead.  Routes all
   multi-token prefill through XMX-accelerated GEMM.

Production data confirms Q patterns: 1-8 TG, 32-127 cache reuse,
128+ full reprocess.  At 32K F16/BF16 FA-on: 356 -> 670 t/s.

* ggml-sycl: fix F16 cache + MKL FA multi-turn corruption; add gate guards

Two changes:

1. Always copy F16 K/V to dense row-major buffers before MKL GEMM.
   Previously F16 was read in-place with raw tensor strides. During
   multi-turn conversations, the accumulated KV cache had different
   stride properties than a fresh prefill, producing corrupted outputs.
   Now dense F16 gets a fast memcpy; interleaved (Gemma) gets a strided
   copy kernel. This matches what the quantized paths already did through
   to_fp16_sycl.

2. Gate MKL FA on unsupported op params (max_bias, logit_softcap, batch
   dim mismatch) and pathological F16 strides (nb[1] not a multiple of
   ne[0]*2). These conditions would previously crash inside the MKL
   kernel. Pathological strides (test-only) and ALiBi/softcap fall
   through to TILE/VEC which handle them correctly.

The stride check uses modulo rather than equality, so both dense
(nb1 == ne0*2) and interleaved (nb1 == H * ne0*2) pass — all real
models use these layouts. Only test cases with overlapping rows
(nb1=32 or nb1=75 for ne0=40) are blocked.

Thanks to hmscider for the oneDNN FA PR (ggml-org#25222) which surfaced the
same insight: always normalize inputs to contiguous F16 before GEMM.

Co-Authored-By: Claude Code using DeepSeek-V4-Pro <noreply@anthropic.com>

* fattn-mkl: fix quant+GQA KV strides, tighten MKL gate, add K>=1024 tests

Adding K>=1024 flash-attn test cases surfaced several MKL bugs:

- Quant K/V with a padded seq-view (real KV cache) used the wrong
  strides in the dequant path... only the true Gemma interleave
  layout should reconstruct strides. nb[2] vs ne[1]*nb[1]
- Gate was firing on shapes the kernel doesn't handle: head_dim < 64
  or not a multiple of 64, MHA, attention sinks, and
  bf16 decode... fell through to vec which no bf16 case.

Gate MKL to the validated envelope: gqa>=2, head_dim 64 through 512
(has to be a multiple of 64) with matching K/V head size, mask,
no sinks/alibi/softcap... everything else falls back to tile.
Covers Qwen Dense/MoE and Gemma4 Dense/MoE

Ran test-backend-ops -o FLASH_ATTN_EXT: 3641/3641 pass.
Perplexity unchanged... 6.7267 MKL vs 6.7290 stock using
Qwen 27b q5_k_xl

* Update ggml/src/ggml-sycl/fattn.cpp

Co-authored-by: Neo Zhang <zhang.jianyu@outlook.com>

* Update ggml/src/ggml-sycl/fattn.cpp

Co-authored-by: Neo Zhang <zhang.jianyu@outlook.com>

* Update ggml/src/ggml-sycl/fattn.cpp

Co-authored-by: Neo Zhang <zhang.jianyu@outlook.com>

* fattn-mkl: bound attention scratch so it doesn't grow with batch or context... also dropped the bf16 comment in fattn.cpp per arthw review.

* Update ggml/src/ggml-sycl/fattn-mkl.cpp

Co-authored-by: Neo Zhang <zhang.jianyu@outlook.com>

* Update ggml/src/ggml-sycl/fattn-mkl.cpp

Co-authored-by: Neo Zhang <zhang.jianyu@outlook.com>

* apply arthw suggestions: enum for dequant modes, macro for wg_size, env-var one-liners

---------

Co-authored-by: Claude Code using DeepSeek-V4-Pro <noreply@anthropic.com>
Co-authored-by: Neo Zhang <zhang.jianyu@outlook.com>
@arthw

arthw commented Aug 3, 2026

Copy link
Copy Markdown
Contributor

In code: commit 11924d4 (tag: b10223, origin/master, master)

It will get performance increase in more LLMs event with fp32 building on B60:

Test Script

  • './build/bin/llama-bench --device SYCL0 -fa 1 -p 4096 -n 0 -m model.gguf'

Environment Configurations

  • env01
    • export GGML_SYCL_ENABLE_MKL_FA=1
  • env02
    • export GGML_SYCL_ENABLE_MKL_FA=0

Per Metric

model metric GGML_SYCL_ENABLE_MKL_FA=0 GGML_SYCL_ENABLE_MKL_FA=1
gemma-4-12b-it-Q5_K_M.gguf pp4096, fa=1, dev=SYCL0, backend=SYCL, ngl=-1 172.73 372.13 (+115.44%)
qwen2-7b-instruct-q4_k_m.gguf pp4096, fa=1, dev=SYCL0, backend=SYCL, ngl=-1 533.98 686.62 (+28.59%)
Qwen3-4B-Q4_K_M.gguf pp4096, fa=1, dev=SYCL0, backend=SYCL, ngl=-1 527.50 617.69 (+17.10%)
Qwen3-8B-Q6_K.gguf pp4096, fa=1, dev=SYCL0, backend=SYCL, ngl=-1 422.50 478.94 (+13.36%)
DeepSeek-R1-Distill-Llama-8B-Q4_0.gguf pp4096, fa=1, dev=SYCL0, backend=SYCL, ngl=-1 458.13 517.34 (+12.92%)
Olmo-3-7B-Instruct-Q4_K_M.gguf pp4096, fa=1, dev=SYCL0, backend=SYCL, ngl=-1 460.75 460.86 (+0.02%)
deepseek-moe-16b-chat.Q4_K_M.gguf pp4096, fa=1, dev=SYCL0, backend=SYCL, ngl=-1 587.67 587.37 (-0.05%)
gpt-oss-20b-Q4_0.gguf pp4096, fa=1, dev=SYCL0, backend=SYCL, ngl=-1 681.75 679.68 (-0.30%)
Bonsai-1.7B-Q1_0.gguf pp4096, fa=1, dev=SYCL0, backend=SYCL, ngl=-1 1290.00 1024.29 (-20.60%)
granite-3.0-3b-a800m-instruct-Q4_0.gguf pp4096, fa=1, dev=SYCL0, backend=SYCL, ngl=-1 1161.84 758.26 (-34.74%)

ggerganov pushed a commit that referenced this pull request Aug 11, 2026
* test new flash_attn test

* rebase and fix to disable subgrou matrices when max_kv_tile == 0

* delete log output

* Add i32 support to cpy and enables the all ops test

* restore the non target ci tests

* comment out of TODO of build-cpu.yml

* fix format
satindergrewal pushed a commit to satindergrewal/llama.cpp that referenced this pull request Aug 12, 2026
ggml-org#25025)

* SYCL: add oneMKL GEMM flash attention for XMX-accelerated prompt processing

* fattn-mkl: fix interleaved dst layout in normalize kernel

- Fix mkl_fa_normalize_head: use interleaved dst layout
  ((query * n_q_heads + head) * DV) matching TILE's
  flash_attn_combine_results. Previously used dense head-major
  layout which wrote head outputs to wrong addresses, corrupting
  attention for all models except Qwen3.6-27B (where GQA=6 heads
  were sparse enough to avoid visible overlap).

- Remove 7 redundant stream->wait() calls — SYCL in-order queue
  already serializes pure SYCL kernel dependencies. Retain only
  the 4 MKL GEMM ↔ SYCL handshake barriers (oneMKL GEMM uses its
  own internal queue that does not respect SYCL in-order).

- Remove unused dst_row_stride, diagnostic clutter, and dead
  K/V hex dump (fa_diag block in fattn-mkl.cpp).

- Add MKL_FA_DISABLE=1 env var for A/B testing.
- Add FA-DISP watchdog (MKL_FA_DEBUG=1) and FA-DIAG output
  fingerprint (MKL_FA_DIAG=1) in fattn.cpp.

Tested: Gemma-4-26B, Gemma-4-31B, Qwen3.6-27B, Qwen3.6-35B-A3B
Perf (B70/Battlemage, 32K, q8_0 KV):
  Gemma-4-26B:  1473 t/s MKL vs 746 TILE (1.97x)
  Qwen3.6-27B:   609 t/s MKL vs 330 TILE (1.85x)

Co-Authored-By: Claude Code on DeepSeek-v4-Pro

* Thank you for the review feedback: rename env vars, use GGML_LOG_INFO, document in SYCL.md

Completed the following:
- Rename MKL_FA_DISABLE → GGML_SYCL_ENABLE_MKL_FA (inverted: 0 to disable)
- Rename MKL_FA_DEBUG → GGML_SYCL_MKL_FA_DEBUG
- Rename MKL_FA_DIAG → GGML_SYCL_MKL_FA_DIAG
- Replace fprintf(stderr, ...) / fflush(stderr) with GGML_LOG_INFO() macro
- Document all three env vars in docs/backend/SYCL.md under Runtime
- Add comment explaining MKL FA activation trigger (flash-attn + quantized
  KV cache + batch-size >= 1024 + n_kv >= 1024)

Resolves review feedback from arthw.
Again, thank you!!!

Co-Authored-By: Claude Code on DeepSeek-v4-Pro

* Thank you for the review feedback round 2: use ggml_sycl_get_env, remove dup waits, gate perf macros

- Replace raw getenv() with ggml_sycl_get_env() in all 4 env-var checks
  (fattn.cpp: GGML_SYCL_ENABLE_MKL_FA, GGML_SYCL_MKL_FA_DEBUG,
   GGML_SYCL_MKL_FA_DIAG; fattn-mkl.cpp: GGML_SYCL_MKL_FA_DEBUG)
- Remove duplicated stream->wait() before ev.wait_and_throw() in GEMM
  KQ and GEMM VKQ — ev.wait_and_throw() already waits for completion
- Gate MKL_ACCUM macro behind do_print so timing accumulators are
  no-ops in normal operation
- Remove redundant MIT/Intel copyright header from fattn-mkl.cpp
- Remove unused #include <cfloat>
- Expand SYCL.md MKL FA docs with step-by-step activation trigger
  and example llama-cli command

Again, thank you!!!

Co-Authored-By: Claude Code on DeepSeek-v4-Pro

* fattn-mkl: enable MKL FA for all KV cache types

Remove the quantized-only restriction on MKL activation — the MKL
kernel converts any non-F16 K/V to F16 via to_fp16_sycl before GEMM,
so F16 (default), BF16, and F32 caches all benefit from XMX hardware
acceleration.  The type restriction was an unnecessary gate.

Before (F16/BF16 default cache + FA on at 32K prefill): ~356 t/s (TILE path)
After:  ~670 t/s (MKL path, matching quantized-cache baseline)

Minimal change: two conditions removed, one comment updated in fattn.cpp.
No kernel or conversion code changes — the dequant pipeline already
covers all types.

* fattn-mkl: rename mkl_disable -> mkl_enable for clarity

* fattn-mkl: refine MKL FA dispatch gates

Three changes:
1. Remove quantized-only restriction - MKL FA activates for all
   KV cache types (F16 default, BF16, F32, quantized).  The MKL
   kernel converts non-F16 K/V via to_fp16_sycl before GEMM.
2. Rename mkl_disable -> mkl_enable to match env var
   (GGML_SYCL_ENABLE_MKL_FA).
3. Replace batch-size threshold with Q->ne[1] >= 32 gate.
   Keeps TG (Q=1) and MTP drafts (Q=3-8) on VEC path where
   fused kernel beats MKL launch overhead.  Routes all
   multi-token prefill through XMX-accelerated GEMM.

Production data confirms Q patterns: 1-8 TG, 32-127 cache reuse,
128+ full reprocess.  At 32K F16/BF16 FA-on: 356 -> 670 t/s.

* ggml-sycl: fix F16 cache + MKL FA multi-turn corruption; add gate guards

Two changes:

1. Always copy F16 K/V to dense row-major buffers before MKL GEMM.
   Previously F16 was read in-place with raw tensor strides. During
   multi-turn conversations, the accumulated KV cache had different
   stride properties than a fresh prefill, producing corrupted outputs.
   Now dense F16 gets a fast memcpy; interleaved (Gemma) gets a strided
   copy kernel. This matches what the quantized paths already did through
   to_fp16_sycl.

2. Gate MKL FA on unsupported op params (max_bias, logit_softcap, batch
   dim mismatch) and pathological F16 strides (nb[1] not a multiple of
   ne[0]*2). These conditions would previously crash inside the MKL
   kernel. Pathological strides (test-only) and ALiBi/softcap fall
   through to TILE/VEC which handle them correctly.

The stride check uses modulo rather than equality, so both dense
(nb1 == ne0*2) and interleaved (nb1 == H * ne0*2) pass — all real
models use these layouts. Only test cases with overlapping rows
(nb1=32 or nb1=75 for ne0=40) are blocked.

Thanks to hmscider for the oneDNN FA PR (ggml-org#25222) which surfaced the
same insight: always normalize inputs to contiguous F16 before GEMM.

Co-Authored-By: Claude Code using DeepSeek-V4-Pro <noreply@anthropic.com>

* fattn-mkl: fix quant+GQA KV strides, tighten MKL gate, add K>=1024 tests

Adding K>=1024 flash-attn test cases surfaced several MKL bugs:

- Quant K/V with a padded seq-view (real KV cache) used the wrong
  strides in the dequant path... only the true Gemma interleave
  layout should reconstruct strides. nb[2] vs ne[1]*nb[1]
- Gate was firing on shapes the kernel doesn't handle: head_dim < 64
  or not a multiple of 64, MHA, attention sinks, and
  bf16 decode... fell through to vec which no bf16 case.

Gate MKL to the validated envelope: gqa>=2, head_dim 64 through 512
(has to be a multiple of 64) with matching K/V head size, mask,
no sinks/alibi/softcap... everything else falls back to tile.
Covers Qwen Dense/MoE and Gemma4 Dense/MoE

Ran test-backend-ops -o FLASH_ATTN_EXT: 3641/3641 pass.
Perplexity unchanged... 6.7267 MKL vs 6.7290 stock using
Qwen 27b q5_k_xl

* Update ggml/src/ggml-sycl/fattn.cpp

Co-authored-by: Neo Zhang <zhang.jianyu@outlook.com>

* Update ggml/src/ggml-sycl/fattn.cpp

Co-authored-by: Neo Zhang <zhang.jianyu@outlook.com>

* Update ggml/src/ggml-sycl/fattn.cpp

Co-authored-by: Neo Zhang <zhang.jianyu@outlook.com>

* fattn-mkl: bound attention scratch so it doesn't grow with batch or context... also dropped the bf16 comment in fattn.cpp per arthw review.

* Update ggml/src/ggml-sycl/fattn-mkl.cpp

Co-authored-by: Neo Zhang <zhang.jianyu@outlook.com>

* Update ggml/src/ggml-sycl/fattn-mkl.cpp

Co-authored-by: Neo Zhang <zhang.jianyu@outlook.com>

* apply arthw suggestions: enum for dequant modes, macro for wg_size, env-var one-liners

---------

Co-authored-by: Claude Code using DeepSeek-V4-Pro <noreply@anthropic.com>
Co-authored-by: Neo Zhang <zhang.jianyu@outlook.com>
zoq pushed a commit to gagallo7/qvac-fabric-llm.cpp that referenced this pull request Aug 12, 2026
…ml-org#26566)

* test new flash_attn test

* rebase and fix to disable subgrou matrices when max_kv_tile == 0

* delete log output

* Add i32 support to cpy and enables the all ops test

* restore the non target ci tests

* comment out of TODO of build-cpu.yml

* fix format
huaxel pushed a commit to huaxel/CachyLLama that referenced this pull request Aug 12, 2026
…ml-org#26566)

* test new flash_attn test

* rebase and fix to disable subgrou matrices when max_kv_tile == 0

* delete log output

* Add i32 support to cpy and enables the all ops test

* restore the non target ci tests

* comment out of TODO of build-cpu.yml

* fix format
mndodd added a commit to mndodd/llama.cpp that referenced this pull request Aug 12, 2026
43 upstream commits, 15 of them in our paths. Four conflicts, resolved as follows.
The dangerous change in this range did NOT conflict -- see (2).

1. ggml/src/ggml-sycl/element_wise.cpp -- TOOK UPSTREAM VERBATIM.
   ggml-org#25946 landed upstream as 11b068d. We had been carrying it as a cherry-pick of
   the then-unmerged PR (ca7c42a) plus two commits of our own stacked on top:
     27d821e  fastdiv for the strided unary index reconstruction
     0595878  fastdiv for the fused-GLU index reconstruction
   Upstream's landed form contains BOTH optimisations by the same mechanism
   (init_fastdiv_values host-side + fast_div_modulo in-kernel, on the strided unary
   path and on all five gated_op_fused_* kernels). All three of ours are therefore
   superseded and are dropped; the file is now byte-identical to origin/master.
   Only semantic difference we give up: ours guarded k > u32 with an exact int64
   fallback, upstream asserts ggml_nelements(dst) < 2^31 instead -- stricter by 2x,
   and unreachable for a GLU activation (~8 GB at f32).
   Our ne>0?ne:1 divisor guard is also dropped; init_fastdiv_values asserts d != 0
   and a ggml tensor always has ne[i] >= 1, so it was defensive, not load-bearing.

2. ggml/src/ggml-sycl/fattn.cpp -- PRECEDENCE PRESERVED, both kernels kept.
   Upstream ggml-org#25025 adds a oneMKL GEMM flash-attention path and gives it
   BEST_FATTN_KERNEL_MKL = 300 -- the value we already use for BEST_FATTN_KERNEL_MMA.
   git flagged the enum collision. It did NOT flag the consequential half: upstream
   places the MKL gate ABOVE our MMA check, and that hunk auto-merged clean.
   MKL's gate is default-ON (GGML_SYCL_ENABLE_MKL_FA=1) and its envelope -- gqa_ratio
   >= 2, head_dim % 64 in [64,512], Q->ne[1] >= 32, K->ne[1] >= 1024, no sinks /
   ALiBi / softcap, with a quantized K/V SKIPPING the F16 stride test -- matches our
   deploy prefill exactly. Taken verbatim it would have silently replaced the
   measured MMA kernel with an unmeasured one, staged the whole q8_0 KV cache to F16
   first, and (per upstream's own note) broken SYCL graph capture replay.
   Resolution: MKL renumbered to 400 so both kernels stay reachable, and its gate
   takes an added !ggml_sycl_fattn_mma_supported(dst) conjunct. MMA wins where MMA
   is supported; MKL keeps its FULL envelope for everything MMA declines, which is
   upstream's intent in every case that is not ours. This is a precedence choice,
   not a revert -- and it is A/B-able without a rebuild:
     GGML_SYCL_FATTN_MMA=0     -> MMA declines, MKL takes the path
     GGML_SYCL_ENABLE_MKL_FA=0 -> MKL off entirely
   Also merged both sides' env-gated instruments, hoisting the kernel selection to a
   single call: upstream re-derived it three times (watchdog, switch, fingerprint),
   so an instrument could disagree with what actually ran. All three now read one
   hoisted `k`. Fixed a latent lie in our own FATTN_DEBUG printer while there --
   BEST_FATTN_KERNEL_ONEDNN was printing as "NONE"; ONEDNN and MKL now print.

3. ggml/src/ggml-sycl/cpy.cpp -- kept ours. Upstream's side of the hunk was empty;
   ggml-org#26005 touched adjacent lines. Our GGML_SYCL_CPY_CENSUS instrument is unchanged.

4. tests/test-backend-ops.cpp -- kept both sides, additive and disjoint (same
   resolution as the 07-28 sync). Ours = the finding-85 deployed-shape MUL_MAT
   sweep; upstream's = m==1 either side of MMVF_MAX_BATCH_SIZE.

Gates run before this commit:
  - cmake-option-audit.sh e9fa078 origin/master -> 7/7 watched options unchanged,
    and --selftest fires (rc=3), so the check is proven able to go red.
  - post-configure asserts: GGML_SYCL / _F16 / _DNN / _GRAPH all ON.

NOT done in this commit, and required before any number from this tree is
comparable to a pre-sync one: rebuild + re-baseline. Upstream changed code under
every arm; ratios within one arm survive, absolutes do not.
brittlewis12 pushed a commit to brittlewis12/llama.cpp that referenced this pull request Aug 17, 2026
…ml-org#26566)

* test new flash_attn test

* rebase and fix to disable subgrou matrices when max_kv_tile == 0

* delete log output

* Add i32 support to cpy and enables the all ops test

* restore the non target ci tests

* comment out of TODO of build-cpu.yml

* fix format
gagallo7 pushed a commit to gagallo7/qvac-fabric-llm.cpp that referenced this pull request Aug 21, 2026
…ml-org#26566)

* test new flash_attn test

* rebase and fix to disable subgrou matrices when max_kv_tile == 0

* delete log output

* Add i32 support to cpy and enables the all ops test

* restore the non target ci tests

* comment out of TODO of build-cpu.yml

* fix format
ravel7524 pushed a commit to ravel7524/llama.cpp that referenced this pull request Aug 30, 2026
…ml-org#26566)

* test new flash_attn test

* rebase and fix to disable subgrou matrices when max_kv_tile == 0

* delete log output

* Add i32 support to cpy and enables the all ops test

* restore the non target ci tests

* comment out of TODO of build-cpu.yml

* fix format
gianni-cor added a commit to tetherto/qvac-fabric-llm.cpp that referenced this pull request Sep 4, 2026
* fix: QVAC-21320 tiled NORM dispatch — never exceed maxComputeWorkGroupCount

Review fix (PR #174): the fused-norm change dispatched GGML_OP_NORM as a direct {ne01, ne02, ne03} grid; on large row counts ne01 can exceed maxComputeWorkGroupCount[0] (spec minimum 65535) and trip the GGML_ASSERT in ggml_vk_dispatch_pipeline, where the previous flattened/tiled dispatch handled arbitrary ggml_nrows.

Restore the flattened {512, 512, N} row tiling on the host (same group as SOFT_MAX/SUM_ROWS) and reconstruct {row, channel, sample} in norm.comp from the flat workgroup id (formula shared with soft_max.comp), with a workgroup-uniform bounds return for the tiling round-up. dst offset is unchanged: flat_row == (samp*nchannels + channel)*nrows + row by construction. No behavioural change for in-range shapes; the fusion's dispatch-count reduction is untouched.

Co-Authored-By: Claude Opus 4.8 (1M context) <noreply@anthropic.com>
(cherry picked from commit 68e6c529ccade61db2f4b0429f0d729a25e49058)

* fix: QVAC-21914 downgrade non-coopmat clip FA hard-disable to AUTO (budget-aware)

The warmup-time hard-disable for GPU projectors without efficient
(coopmat) flash attention replaced AUTO/ENABLED with DISABLED, which
short-circuited the budget-aware AUTO heuristic in
clip_resolve_flash_attn_type(). At high n_pos (image_tile_mode=disabled
with image_max_tokens=4096 -> 16384 ViT patches) the forced explicit
attention path materializes an O(n^2 * n_head) score matrix, growing RSS
to ~12 GB and getting the process lmkd-killed on Pixel 9 Pro
(runQwen35ImageTileModeTokensTest).

Downgrade to AUTO instead and record the inefficiency in
clip_ctx::fa_backend_inefficient, which now also enables the AUTO cutoff
default (previously Mali-detection only) so any non-coopmat backend gets
the per-image budget decision: explicit attention below the cutoff (fast
on scalar-FA GPUs), memory-frugal scalar FA at/above it or when the
explicit scratch would not fit device memory. Explicit user DISABLED is
still honored, and MTMD_CLIP_AUTO_FA_MIN_KV still overrides the cutoff.

(cherry picked from commit cceee2288c686df79634bcb029de629542e45756)

* fix: QVAC-21914 bound ggml-opencl submissions (periodic clFlush + FA q-chunking)

On Galaxy S25 Ultra (Adreno 830, OpenCL) the monolithic 16384-patch ViT
encode (image_tile_mode=disabled, image_max_tokens=4096) faults the GPU
near the end of the encode (Adreno-GSL log_gpu_snapshot fires before any
decode work reaches the device), after which the driver aborts the
process from cl_a8x_cmdbuf_mgr_submit_ibs (os_exit) on the next
submission. Two unbounded behaviours plausibly drive the fault and both
are bounded here:

- ggml_backend_opencl_graph_compute enqueued entire graphs (thousands of
  nodes, ~48 s of GPU work for the failing encode) with no intra-graph
  flush. Now clFlush every GGML_OPENCL_FLUSH_INTERVAL nodes (default 64,
  0 disables) so the GSL command-buffer manager receives bounded
  batches. clFlush submits without stalling the host.

- ggml_cl_flash_attn issued one dispatch covering all q rows; at
  n_q = n_kv = 16384 every workgroup loops the full KV, making a single
  very long kernel. Now chunked along q rows at GGML_OPENCL_FA_MAX_NQ
  rows per dispatch (default 4096, 0 disables) with a clFlush between
  chunks. The split is exact: the kernel resolves its q row relative to
  the Q/O/mask base offsets, is_causal is always 0 (masking is explicit)
  and alibi/sinks depend only on the head index, so shifting the row
  base via byte offsets while shrinking n_q is mathematically identical.
  No .cl kernel changes.

The 512-token image-chunk decode is not implicated: the S25 VLM
benchmark ran 304 full-512-row ubatch decodes cleanly. Only the giant
monolithic encode (5-8x beyond anything previously run on this backend)
triggers the fault.

(cherry picked from commit 82a26df6e5ce24053101a148c0423d54515226bd)

* fix: QVAC-21914 review fixes — work-budget flush, memory-clamp rework, tests

Addresses the pre-merge review findings on the two QVAC-21914 crash-fix
commits (P1/P2 performance, C1/C2 correctness, S1/S2 robustness, K nits):

- ggml-opencl: gate the periodic graph flush on accumulated estimated WORK
  (GGML_OPENCL_FLUSH_WORK_MB, default 512 MB) instead of a bare node
  counter. Per-token LLM decode graphs never reach the budget by
  construction, so the decode hot path stays submission-free; the
  16k-patch encode still flushes dozens of times. Single touch point in
  graph_compute (no more per-fusion-branch duplication).
- ggml-opencl: both tunables move onto ggml_backend_opencl_context,
  resolved once at init with strtol-based parsing (clamp, warn on garbage
  instead of silently disabling the mitigation) and GGML_LOG_INFO'd like
  the file's other env knobs. FA chunking reads the context field.
- ggml-opencl: GGML_ASSERT(is_causal == 0) before the FA chunk loop — the
  kernel's causal-boundary formula needs the TOTAL n_q, so chunks after
  the first would silently corrupt output if causal FA were ever enabled
  here; keep the invariant loud. Explicit n_q == 0 guard.
- clip: rework the AUTO cutoff memory clamp. Total memory now provides the
  STABLE fast-path clamp (explicit scratch <= total/4); free memory (a
  volatile, load-dependent number) may only lower the cutoff further via
  the hard-fit requirement (scratch <= free), never the old free/2
  heuristic that silently pushed normal-size Mali images onto the ~2.6x
  slower scalar-FA path under momentary memory pressure. No memory info at
  all now fails SAFE at a conservative 2048-patch cap instead of trusting
  the raw 4096 default (~3.2 GB scratch at n_head=16). The arithmetic is
  extracted into clip_fa_effective_min_kv() (pure, exposed via clip.h for
  tests).
- tests: test-clip-fa-cutoff (pure CPU, locks in the fast path, the P2
  regression guard, the fail-safe cap and edge cases; passing) and
  test-opencl-fa-chunking (chunked-vs-CPU numerical parity over unchunked
  / exact-chunk / partial-last-chunk / n_q==1, masked and unmasked, with
  GGML_OPENCL_FA_MAX_NQ=64 and a 1 MB flush budget; self-skips without a
  capable OpenCL device — PoCL lacks FP16, so it executes on Adreno-class
  hardware).
- clip warmup comment: note ggml-opencl also lands in the "no efficient-FA
  query" bucket and its giant-encode fault is handled by the submission
  bounding inside that backend.

GGML_OPENCL_FLUSH_INTERVAL (node-count knob) is replaced by
GGML_OPENCL_FLUSH_WORK_MB; GGML_OPENCL_FA_MAX_NQ semantics unchanged.

(cherry picked from commit fc09f36b4435a23b1e261e20ca9d6122d68c2404)

---

b10297 rebase:

- Squash a68b35970: test-opencl-fa-chunking: call
  ggml_backend_load_all() and select the device by backend registry
  name ("OpenCL") instead of substring-matching the device name.

Squashed-with: 649de77eb, a68b35970

* fix: QVAC-21914 make clip_fa_effective_min_kv inline (Windows DLL link)

The pure AUTO-budget helper was defined out-of-line in clip.cpp and
declared in the internal clip.h. On Windows mtmd builds as a shared
library exporting only the MTMD_API-decorated public API; the internal
clip_* symbols are absent from mtmd.lib, so test-clip-fa-cutoff (the
first cross-DLL-boundary consumer of a clip_* symbol) failed to link
(LNK2019). Linux/macOS export all default-visibility symbols, so it
linked there.

Move the function inline into clip.h (with its NO_MEMINFO_CAP constant);
the test and clip.cpp both compile their own copy — no DLL export of an
internal helper. CLIP_AUTO_FA_MIN_KV_MALI_DEFAULT stays in clip.cpp (its
only user). Verified: mtmd + test-clip-fa-cutoff build and the test
passes.

(cherry picked from commit 99d6042207cb255a8334d97546957fbab19bfc66)

* fix: QVAC-21914 address PR review — warn on env clamp, cover n_head guard

- parse_env_i64: GGML_LOG_WARN when an in-range-but-too-large
  GGML_OPENCL_FLUSH_WORK_MB / GGML_OPENCL_FA_MAX_NQ is clamped to max,
  matching the file's convention of logging every overridden value
  (previously the clamp was silent).
- test-clip-fa-cutoff: the n_head=0 case passed total_mem==free_mem==0,
  which short-circuits to the NO_MEMINFO cap before any sqrt(.../n_head)
  branch runs — the div-by-zero guard was never exercised. Pass 16 GB
  total so the total-memory clamp runs with n_head=0; without the guard
  the (int)sqrt(x/0) path would now fail the assertion.

Both from yingying0906's review; Windows/CPU-only surface, no Android
behavior change.

(cherry picked from commit 3da3e05fc459778f343da6ce0b796766544546a7)

* fix: QVAC-21914 flush after enqueuing the budget-crossing node

The graph_compute work-budget flush ran BEFORE the current node was
dispatched: it accounted the node's work, and on crossing the budget
flushed (submitting only the prior batch) then reset the counter to 0 —
so the crossing node started a fresh batch. A large op, or the last
large segment of the graph, could therefore begin an unflushed batch and
be submitted unbounded at the implicit end-of-graph finish, defeating
the bound.

Move the budget check below the dispatch (convert the fused-op
continue chain to if/else so every path reaches one touch point), so the
node that crosses the budget is part of the flushed batch. Reported by
@gianni-cor.

(cherry picked from commit 7056f4cadb60aa255a333314a45c6d520ef88637)

* chore: QVAC-21914 address PR review — flush-cadence wording, nits

Non-behavioral cleanups from the PR #181 review pass:

- ggml-opencl graph_compute: correct the flush-cadence comment. The old
  "per-token decode hot path submission-free by construction" claim was
  false for multi-GB models — a decode step streams the whole model, so
  its estimated work crosses the default 512 MB budget a few times per
  token. Reworded to state that accurately (cost is negligible in
  practice since clFlush is non-blocking, and it is tunable/zeroable to
  make decode fully submission-free). Mechanism unchanged.
- ggml_cl_flash_attn: GGML_ASSERT(q->ne[1] <= INT32_MAX) before the int
  n_q truncation, since the q-chunk loop accumulates into an int and
  derives cl_ulong offsets from it (defensive; not reachable with real
  shapes).
- Consistency: normalize the ticket tag to bare `QVAC-21914` (drop the
  `qvac ` prefix) in clip.h and tests/CMakeLists.txt, matching the .cpp
  files and the fork's QVAC-21257 precedent.
- tests/CMakeLists.txt: move the unconditional test-opencl-fa-chunking
  registration up beside test-copy-tbq-subgroups (its self-skipping
  sibling) instead of sitting right after the LLAMA_MTMD endif() where it
  read as MTMD-gated; add a comment noting it deliberately does not link
  mtmd.

No functional change to the fix; local mtmd + both tests build,
test-clip-fa-cutoff passes.

(cherry picked from commit 2b927cf3910c07c8ccebae984fc8eabc2ffa17b0)

* vulkan: restore the nb00 element stride in the fused norm shader so non-contiguous (permuted) inputs read the right columns

Signed-off-by: Marcus Edel <marcus.edel@collabora.com>
(cherry picked from commit 6053e42d473e0a828682fd282b11a8a5cc198603)

* cuda: Fix OUT_PROD op support claim

Only claim OUT_PROD support for src types ggml_get_to_fp32_cuda can
requantize.

ggml_cuda_out_prod converts non-F32 srcs to F32 before the f32-only
cuBLAS GEMM and aborts on e.g. TQ2_0 which has no CUDA dequantizer.

(cherry picked from commit d52948582510fc8ed168998473f901698485e76d)

* ggml-opencl: build flash-attn kernels without finite-math

The OpenCL kernels are compiled with -cl-finite-math-only and
-cl-fast-relaxed-math, which let the compiler assume no Inf/NaN. The
flash-attention online softmax initialises its running max to -INFINITY
and masks padded scores with -INFINITY, so finite-math miscompiles the
init/masking path.

Compile the flash-attention programs with a relaxed option set that
drops -cl-fast-relaxed-math, -cl-finite-math-only and
-cl-unsafe-math-optimizations (keeping -cl-mad-enable for speed) so the
-inf sentinels behave correctly.

Also harden the strip: erase every occurrence of each flag (not just the
first) and GGML_ASSERT that no finite-math/fast-math/unsafe-math flag
survived, so a future compile_opts spelling/spacing change fails loudly at
load time instead of silently reintroducing the -INFINITY miscompile.

Re-ported onto b9840's rewritten OpenCL flash-attn (upstream PR #14987 +
follow-ups): the original per-dim kernel-compile loop is gone, so the strip
is applied once in ggml_opencl_fa_compile_opts(), the single site every FA
variant (F16/F32/F32_F16/Q8_0/Q4_0/PRE and _SPLIT) is compiled through.
Squashed re-port of 0cbe36259 + c1dace72b (finite-math part).

(cherry picked from commit 348a910361bdd8b4bd0f8e70c8e793ef7fbc5ee5)

* ggml-opencl: treat null attention mask as bidirectional, not causal

The flash-attention dispatch inferred causal masking from shape with
`is_causal = (mask == NULL && n_q > 1 && n_q == n_kv)`. A null mask means
no masking, i.e. bidirectional attention (the SigLIP vision and embedding
encoders), while causal attention always supplies an explicit causal mask
in this codebase (llama-graph.cpp build_attn passes a kq_mask filled with
-INFINITY). The heuristic therefore wrongly made the bidirectional
Qwen3-VL vision tower attend causally, so each patch only saw earlier
patches and the image embedding was corrupted.

Set is_causal = 0 unconditionally; causality is always expressed via the
explicit mask. This cannot regress the LLM, which already passes a real
causal mask (is_causal was already 0 for it) and relies on that mask.

Document the invariant in ggml_cl_flash_attn: a null mask is treated as
bidirectional, so any caller needing causal masking must supply an explicit
causal mask rather than relying on shape inference.

Re-ported onto b9840's rewritten OpenCL flash-attn; b9840's own q-chunking
path already GGML_ASSERTs is_causal == 0, so this is consistent with the
existing code. Squashed re-port of 51dbb1756 + c1dace72b (is_causal part).

(cherry picked from commit 7ae4bc939f21f3b5c72ab3caa01c806605d6a15b)

* ggml-opencl: add trailing barrier in f32/f16 flash-attn tile loop + guard upscale zero dims

The f32/f16 flash-attention kernels load K/V tiles into local memory,
barrier, read them, then loop to overwrite the tiles for the next K/V block
without a trailing barrier. Out-of-range lanes (the last partial BLOCK_M
block) `continue` past the read and race ahead into the next tile load while
active lanes are still reading l_k/l_v. With n_kv > BLOCK_N (e.g. the
bidirectional vision tower, n_kv=247) this corrupts the shared tiles.

Add a trailing barrier(CLK_LOCAL_MEM_FENCE) at the end of the K/V block loop
and guard the score computation with `if (my_query_row < n_q)` instead of an
early continue. flash_attn_f32_f16.cl already uses that guard + trailing
barrier after b9840's redesign, so it is left unchanged.

Also guard zero source dimensions in ggml_cl_upscale: the sf* scale factors
divide by the source dims, so a zero source dim yields +inf; the existing
early-exit only covered zero destination dims.

Re-ported onto b9840's rewritten OpenCL flash-attn (b9840 widened the score
unroll to j += 4; only the divergence guard + trailing barrier are re-applied,
the body is unchanged). Squashed re-port of dc64397d2 + b7ad6d4e2.

The barrier fix is a GPU-scheduling race whose only proof is on-device; the
original b7ad6d4e2 validated on S25 Ultra / Adreno 830 (Qwen3-VL GPU vision
projector matches CPU exactly, Delta 0.0 pp, ~26% faster on encode). Re-verify
on-device before merge.

(cherry picked from commit ce54dd55ab6dd590ee597c868b12071b11d01732)

---

b10297 rebase:

- Squash 97a1ecd12: flash_attn_f32.cl: apply the divergence guard
  unconditionally and drop the FA_SG<64-only trailing barrier; the
  tile race reproduces even on a single 64-wide Adreno subgroup
  (Adreno 830).

Squashed-with: 278014e04, 97a1ecd12

* server tests: detect cache restore via timings, not exact log text

The upstream target of this former fixup! (50e0ad08f, --clear-idle
#20993) already landed upstream, so this stays a standalone commit.

Relying on exact log text is brittle, especially across rebases with
upstream changes; use the timings fields instead and drain remaining
logs for test cleanliness.

* ci: install jinja2 explicitly in the venv so test-jinja-py doesn't depend on the torch pin surviving pip install

Signed-off-by: Marcus Edel <marcus.edel@collabora.com>
(cherry picked from commit 2d3f4034bd6dcf84089c2874f1d1ab26f654576f)

* ci: add backend op-coverage manifest guard + SVE variant tripwire

A supports_op() regression never fails test-backend-ops: the case falls
back to CPU and is reported 'not supported [backend]', i.e. skipped.
e09ae0b71 removed Q4_1/Q4_K from OpenCL MUL_MAT supports_op with zero
test failures

---

b10297 rebase:

- Squash e8fb282d9: add the UPSCALE f32 bilinear|antialias row to the
  opencl-pocl op-coverage manifest.

Squashed-with: 5587e68af, e8fb282d9

* tests: add no-mask n_q == n_kv FLASH_ATTN_EXT cases (vision-tower shape)

The FA sweep uses kv in {113, 512, 1024} x nb in {1, 3, 32, 75}, so nb
never equals kv and the mask==NULL && n_q==n_kv shape -- exactly what a
ViT self-attention layer produces -- is never exercised. A backend that
infers causality from that shape (OpenCL's is_causal heuristic,
ggml-opencl.cpp:15131) silently computes causal attention for the whole
vision tower and no CI test goes red.

Add explicit bidirectional cases at n_q == n_kv == 247 and 256 for head
sizes 64 and 80 (both in the OpenCL FA supported-dims table). 247 (odd)
additionally leaves partial tiles for any power-of-two tile size; 256 is
the aligned control separating causality bugs from tiling bugs.

Expected red on OpenCL until the is_causal heuristic is removed
(re-port of b9840 7ae4bc939); green on CPU/Vulkan/CUDA/Metal.

Note for landing: order this commit after the is_causal fix so the
series stays bisectable-green on OpenCL hardware.

Assisted-by: Claude (Anthropic AI)

Co-Authored-By: Claude Fable 5 <noreply@anthropic.com>
Claude-Session: https://claude.ai/code/session_01DbaV79zWiPZJq1LrdWGDTj
(cherry picked from commit a0fe46a41d5c8e02f9923049c0a16cbaef8d7439)

* tests: add partial-tile FLASH_ATTN_EXT cases across KV-type kernel variants

n_q=33 with n_kv=513 leaves out-of-range query lanes for any pow2 query
tile (Adreno OpenCL uses BLOCK_M=64 for dk 64/128) and a 1-valid-row
final KV tile (513 = 16*32 + 1), with 17 tile-loop iterations worth of
barrier crossings. Tiled kernels must keep out-of-range lanes inside the
tile loop for the barriers while excluding them from the score loop; an
early continue past the tile barrier (the b9840 ce54dd55a race class) or
a missing trailing barrier corrupts the shared K/V tiles for in-range
lanes.

Cover each KV type separately: backends that specialize kernels per KV
type (OpenCL picks flash_attn_f32.cl for f32 KV, f32_f16(+split) for f16
KV since n_kv=513 crosses the split threshold, and the q8_0/q4_0 tiled
kernels for quant KV) would otherwise leave those variants untested at
this shape. Note flash_attn_f16.cl itself is only reachable with an f16
Q tensor, which test-backend-ops never generates (Q is always f32) -- it
stays covered only by code review.

The race is scheduling-dependent: in-order devices (pocl) and drivers
with cooperative-matrix FA paths will likely pass even with the bug;
these shapes make the sweep able to catch it on the affected hardware
class (original repro: Adreno 830).

Validated: CPU supports all 5 cases; Vulkan RADV 5/5 PASS on 7900 XTX
and Raphael iGPU.

Assisted-by: Claude (Anthropic AI)

Co-Authored-By: Claude Fable 5 <noreply@anthropic.com>
Claude-Session: https://claude.ai/code/session_01DbaV79zWiPZJq1LrdWGDTj
(cherry picked from commit ab6ecf6b4ea23533640988f0557b11c48106bf01)

* vulkan: add GFX1151 q8_0 BK64 matmul path

Add the tuned cooperative-matrix path without experimental runtime controls.

(cherry picked from commit 9675568aeb84d91aa0653bcf93b0e6aa225eaa7d)
Co-authored-by: Guilherme Gallo <guilherme.gallo@collabora.com>

* cpu: optimize DeepSeek4 HC post

Vectorize the fused HC post path for ARM and x86 while preserving scalar fallback behavior.

(cherry picked from commit 7c21b4a5400ac9c99576dbce401bf26b2209175f)

* vulkan: add Lightning Indexer

Assisted-by: GPT-5.6 Sol
(cherry picked from commit 5f41c2e29ee5b5cdcee383fef2208173604c7704)
Build-fix-hoisted-from: d19b9ddbadc9 (drop stray closing brace after ggml_vk_lightning_indexer)

* vulkan: add DeepSeek4 HC comb

Assisted-by: GPT-5.6 Sol
(cherry picked from commit 9afa822b70fa99a08d14f9c3369ab1a674332a55)

* vulkan: add DeepSeek4 HC pre

Assisted-by: GPT-5.6 Sol
(cherry picked from commit e5d7104869fe4a9d4d1aa6d1b5d0e6a5bb963bd0)

* vulkan: add DeepSeek4 HC post

Assisted-by: GPT-5.6 Sol
(cherry picked from commit 16d647b7046d7897eaee7c6f9d4ceba20dcd3f17)

* tests: cover DeepSeek4 fused operations

Add fused-versus-fallback correctness and benchmark cases for HC post and the Lightning Indexer.

Assisted-by: GPT-5.6 Sol
(cherry picked from commit b5dcb61adceb221ec3c688d6c957966a7bf4f5c0)

---

b10297 rebase:

- skip_backend for DSV4_HC_POST_BIT_EXACT matched only the "CUDA" reg
  name, but HIP builds register the same backend as "ROCm"
  (GGML_CUDA_NAME), so the FMA-contraction skip never fired there and
  the bit-exact gate failed deterministically on ROCm (6 cases,
  ERR=1.0).

* vulkan: add typed K cache support to Lightning Indexer

Generate per-K-type generic pipelines and 32/64-head CM1 and CM2 variants, with quant-aware stride handling and guarded pipeline selection.

CM1 stages decoded FP16 tiles under shared-memory limits; CM2 uses FP16 decode callbacks for quantized K tiles.

Assisted-by: GPT-5.6 Sol
(cherry picked from commit d2414ec43e06309824c332c5e00bbb90c0bec5be)

* tests: cover typed Lightning Indexer K caches

Exercise matrix paths for 32- and 64-head layouts across supported K-cache types, including dispatch-tail boundaries and a strided-Q scalar fallback.

(cherry picked from commit 1a8d1a67ebedd5cdd71ea874d339400da50d0449)

* vulkan: support quantized concat

Keep quantized cache concatenation on the GPU to avoid per-layer CPU fallbacks and excessive graph splits.

(cherry picked from commit 7e530286b4d36b2c70ceec445e380bcdbd9459f8)

* metal: support quantized concat

(cherry picked from commit 581bbdafb9e4a101970787107cb9baceb3c2823b)

* metal: support strided f16 add

Keep DeepSeek V4 mask construction on the GPU with a vectorized path for strided inputs.

(cherry picked from commit 807cc6aa0bf6acfaf741be41494643632a059dd7)

* meta: handle Lightning Indexer split state

Treat the fused indexer conservatively like the other DeepSeek4 fused operations instead of aborting in meta backends.

(cherry picked from commit 7aad3a2ef844fce37c729d3329f159105ad52712)

* ci: extend macOS ARM test timeout

Allow the expanded backend operation suite to complete on slower Apple runners.

(cherry picked from commit edbc216ba6cf7b7e2bc482fd07377f6aa2a54e1e)

* ci: make nproc portable and verify it in self-hosted deps

ci/run.sh uses nproc for build -j and quantize thread counts, but macOS
runners often lack GNU coreutils, so $(nproc) silently expanded to
nothing, polluting the logs

Define a sysctl-backed shim when nproc is missing, fail setup outright
when neither is available, and surface the gap in
gg_check_build_requirements

On the Graviton jobs, install coreutils explicitly and assert nproc
resolves at the end of the Dependencies step so a missing tool fails
there instead of mid-run

* ci: fail fast when jinja2 is missing from the CI venv

Assert jinja2 is importable by the same python3 the test spawns and
abort setup with a clear error instead.

* Add CHANGELOG.md

Signed-off-by: makaveli10 <vineet.suryan@collabora.com>

* ggml-opt: abort training when the backend graph compute fails instead of accumulating garbage results

Signed-off-by: Marcus Edel <marcus.edel@collabora.com>

* ggml-metal: default n_cb back to 1 on iOS to fix the A18 GPU hang when resuming finetuning from a pause checkpoint

Signed-off-by: Marcus Edel <marcus.edel@collabora.com>

* ggml-metal: only track memory ranges when the encoder is concurrent, fixing the A18 GPU hang from barriers encoded into serial encoders

Signed-off-by: Marcus Edel <marcus.edel@collabora.com>

* ggml-metal: avoid per-iteration whole-block copies in the quantized OUT_PROD kernels so finetuning backward passes stay under the iOS GPU watchdog

Signed-off-by: Marcus Edel <marcus.edel@collabora.com>

* vulkan: drain glslc stdout/stderr concurrently

This makes shader compile fail fast.

execute_command() drained the child stdout pipe to EOF before touching
stderr. Poll both pipes and read whichever is ready until both hit EOF

A glslc invocation that emits more than a pipe buffer of diagnostics
(e.g. hundreds of errors from a broken shader variant) blocks in
write(2) on the full stderr pipe, the parent blocks in read(2) on
stdout, and the whole shader-gen step deadlocks instead of failing the
build

Assisted-by: Claude Fable 5 <noreply@anthropic.com>

* QVAC-21550 infra: roll out canonical security baseline (TruffleHog + CodeQL) to qvac-fabric-llm.cpp (#177)

* QVAC-21550 infra: add canonical security baseline caller (TruffleHog + CodeQL)

* QVAC-21550 fix: drop paths-exclude (unsupported by reusable security v0)

* QVAC-21550 fix: drop secrets: inherit (baseline needs no repo secrets; github.token suffices)

* QVAC-21550 fix: repin security baseline to qvac-actions 0.2.0 (buildless c-cpp)

(cherry picked from commit 4b81205eecc25e6ae43f1b88b317da646d7606a6)

* QVAC-22747 infra: add weekly schedule to CodeScan caller (#189)

The canonical CodeScan baseline currently runs on push/PR only, so the
QVAC-19056 commitment to a weekly scan is unmet. Add a scheduled cron
(staggered per repo across the fleet) alongside the existing
push / pull_request / workflow_dispatch triggers. No other change.

(cherry picked from commit 5922019a656e1ad84e27b72f237ae49a354c6615)

* QVAC-22740 infra: bump CodeScan caller to qvac-actions 0.3.0 (#195)

Bump the reusable security workflow pin 0.2.0 -> 0.3.0 (SHA bbb0740e). 0.3.0
adds the findings-export artifact (export-report default on), so each run now
publishes a downloadable security-scan-report (findings.json/md + SARIF).

(cherry picked from commit de769869f0e29d4e002c28e8e796ab6cd8e0593f)

* chore(ci): remove qvac-collabora-merge from CODEOWNERS

Per QIP tier-1-approval-change, only management and team-lead teams
should be listed as code owners.

(cherry picked from commit 7b02a3c3c13c271dbf1b1a1059c568ac6cb1b96c)

* squash! metal: port OUT_PROD, SILU_BACK, SOFT_MAX_BACK, RMS_NORM_BACK ops to split architecture

Drop the downstream SILU_BACK port: upstream b10297 ships its own typed kernel_silu_back_<type> pipeline, op encoder, and supports_op case, so the downstream _4-suffix variant became a duplicate definition (redefinition errors in ggml-metal-device.cpp/ggml-metal-ops.cpp and a duplicate case label in ggml-metal-device.m supports_op).

Assisted-by: Claude Fable 5

Co-Authored-By: Claude Fable 5 <noreply@anthropic.com>
Claude-Session: https://claude.ai/code/session_01CBhFiEgFVPcnDBpdP4e6Yk

* tests: relax dsv4 bit-exact check to ulp-based tolerance

Signed-off-by: Marcus Edel <marcus.edel@collabora.com>

* squash! hip: allow-list rwkv_wkv_f32<128> VGPRs

Refresh the mul_mat_q<Q2_K, 64, true> mangled name

Upstream 1a064ab09 (NVFP4 W4A4 activation quantization) added a const
float * y_scale kernel argument, which changed the mangled signature and
orphaned the previous allow-list entry

* ggml-webgpu: fix CI errors from #25025 and #25262 (#26566)

* test new flash_attn test

* rebase and fix to disable subgrou matrices when max_kv_tile == 0

* delete log output

* Add i32 support to cpy and enables the all ops test

* restore the non target ci tests

* comment out of TODO of build-cpu.yml

* fix format

* tests: add more cases and perf mode support for backward GDN

* ggml-metal: add getter function for threadExecutionWidth

* ggml-metal: allow ggml_metal_library_get_pipeline_* fns to derive thread dispatch sizes

* ggml-metal: optimizes gated-delta-net back op

Replicates the same optimizations applied to the
vulkan version of the GDN back kernel

* ggml-cuda: adds support for backwards gated delta-net

* ggml-cuda: adds sigmoid-back

* ggml-cuda: implement geglu backward

Signed-off-by: makaveli10 <vineet.suryan@collabora.com>

* ggml-cuda: implement ssm-conv-back-sx

* ggml-cuda: implement ssm-conv-back-c

* QVAC-23075 feat: add VisionPsy Nano and its Flash preprocessing rule

VisionPsy Nano is a siglip encoder with an idefics3-style pixel-shuffle merge and a
single mm_fc projector, so route it through clip_graph_siglip and reuse the llava-uhd
preprocessing. Four differences from idefics3, all taken from get_image_string() in
the reference processors.py: position embeddings are interpolated to the actual patch
grid, the overview image is emitted first, the delimiters are <|global_image|> and
<row_%d_col_%d> with no fake_token_around_image and no row-end token, and each image
carries an <image: N> ordinal label once a prompt holds more than one. That label is
not in the vocab, so it goes through BPE as text.

When the slice grid is 1x1 the overview and the single slice are the same crop, so
only one of them is sent, and it is the slice. GlobalAndSplitImages.forward returns
the split patch untouched for that grid, before it resizes a global patch, and the two
paths do not share a resize kernel: the slice is rendered with image_resize_algo_rf
(bicubic) and the overview with image_resize_algo_ov (bilinear).

The published mmproj GGUFs declare clip.projector_type = "custom", so add a read-only
alias table mapping that string onto PROJECTOR_TYPE_VISIONPSY. The alias is gated on
general.name as well, so another model shipping "custom" is not silently loaded as
VisionPsy. PROJECTOR_TYPE_NAMES stays canonical and "visionpsy" is what we write out.

lookup_token() returns LLAMA_TOKEN_NULL on a miss and no caller checks it, so a vocab
without <|global_image|> would splice a garbage id into the prompt. It now throws, but
only when there is a vocab to look in: mtmd_get_memory_usage builds a context with no
text model and every token misses there, which otherwise cost every llama-server start
with this mmproj its mmproj VRAM reservation.

The base and Flash checkpoints differ in exactly one thing, resize_to_max_side_len.
Base always stretches the long side to max_img_size; Flash rounds it up to a whole
number of slices and only then caps it, so an image below the cap keeps its own
resolution. Their mmprojs are byte-different but declare identical vision hparams, so
nothing in the GGUF distinguishes them and Flash was silently running base
preprocessing, inflating a 256x256 image into a 4x4 grid of upscaled slices.

Add clip.vision.preproc_no_upscale, read from the GGUF when present, overridable
through mtmd_context_params.image_no_upscale and the new --image-no-upscale flag. The
published mmprojs do not carry the key, so today the flag is what selects the variant;
a re-converted Flash mmproj would need no flag. Tri-state, -1 keeps the model default,
so every existing caller is unaffected.

The short side is computed in double. The reference is math.ceil(short * scale / p) on
Python floats, and where the true quotient lands on an integer, float32 and exact
integer arithmetic each disagree with it by a whole slice row: over every size pair up
to 3000x3000, float32 differs in 171 cases and exact integers in 92, double in none.
image_size == 0 is legal in general but calc_size_no_upscale divides by it, so it is
rejected at load, where the rest of the loader reports bad config, rather than
asserting per image.

Flash on 1024x768 goes from 13 slices to 5, encode 361 ms to 139 ms, prefill 858
tokens to 338; on 256x256 from 17 slices to 1, 474 ms to 38 ms, 1118 tokens to 78.
DocVQA ANLS over 30 examples is unchanged at 0.825 versus 0.824, as expected, since
those scans are larger than the cap and the two rules converge above it.

Both checkpoints are registered in the vision test list. They use the published
GGUFs, which -hf resolves to both the language model and the mmproj, and answer "the
new york times" to the harness default prompt, so they pass the existing
new-york / men-walk assertion. Flash needs the flag on its row, since its published
mmproj carries no clip.vision.preproc_no_upscale key and would otherwise duplicate
the base row's preprocessing.

(cherry picked from commit fc07c46a16298db9514ba66dddb3573eb74d9829)

* QVAC-23075 fix: ask the Adreno transpose predicate about the parent, not the view

enable_adreno_trans_weight gates three things as a unit: the transpose in set_tensor,
the restore in get_tensor and the adreno gemm/gemv selection in mul_mat. set_tensor
returns early for views and always decides on the parent, but get_tensor asked about
the view, so the two could disagree on the shape predicate and pair a transposed
writer with a non-transposed reader.

Unrelated to VisionPsy, and kept separate for that reason. The shape predicate itself
already rejects the q8_0 layouts the transpose asserts on, so this only closes the
view-versus-parent gap.

(cherry picked from commit 7a6425a3add96f2598110a5953786866c9d49e0c)

* QVAC-23075 fix: reconstruct the q8_0 parent before reading a view back on Adreno

The SoA buffers cover the whole parent and the Adreno transposed layout is indexed by the
parent's row count, but get_tensor sized its staging buffer and its dispatch from the view.
A view with fewer rows than its parent therefore strided the transposed source wrongly, and
once the dispatch rounded the row count up to the 64-wide work group it wrote past the end
of the buffer: a 4096x1 view of a 4096x4096 parent allocates about 4.3 KB and the kernel
restores 64 rows into it, about 278.5 KB. Reconstruct the parent instead, exactly as
set_tensor already does, and read the view out of it at view_offs.

The kernel also gets the row guard the rounded-up dispatch needs, which closes the same
overflow for any parent whose row count is not a multiple of 64.

(cherry picked from commit 62162f7fa0bdaf5849bc19f6d62418c8c94205be)

* QVAC-23075 fix: size the aspect-preserving refine in double, not float

calc_size_preserved_ratio scaled in float32 while both reference processors do it in Python
floats. Where the true product lands exactly on a multiple of align_size, float32 rounding
pushes it just past and the ceil buys a whole extra row or column of slices. 960x720 with
image_size 512 and longest_edge 2048 is the common 4:3 case: the reference refines to
2048x1536 and slices 4x3, float32 gave 2048x2048 and sliced 4x4, so the image gained 4
slices and 256 image tokens the model was never trained to see.

Verified against the local Metal build on the base checkpoint: 960x720 drops from 17 image
encodes to 13, that is a 4x4 grid to a 4x3 grid plus the overview. 1024x768, 640x480 and
256x256 are unchanged, their scale factors are exact in float32. Sweeping every 16-pixel
size pair up to 3000x3000, float32 and double disagree on 10 of 33856, all of them 4:3 or
3:4. The idefics3 and pixtral paths share this helper and both references are also double,
so they move the same way.

(cherry picked from commit a151b019c2bf90393a58bb1162d4ae2a54bceb37)

* QVAC-23075 fix: reject a no-upscale cap below one slice

calc_size_no_upscale() clamps both target sides into [image_size, image_longest_edge], and
std::clamp requires lo <= hi, so GGUF metadata with a cap smaller than the slice size is
undefined behaviour rather than merely a bad size. Validation already demanded both be
positive; demand the ordering too, next to it and after the flag override, so it covers the
CLI flag as well as the GGUF key.

(cherry picked from commit 0b61b39c5698b69419a7f535c518a5c4e2eb927c)

* QVAC-23075 fix: keep the NUL terminator out of string_format's result

clip-impl.h's string_format returned std::string(buf.data(), buf.size()) over a buffer sized
size + 1, so the terminating NUL sat inside the string and length() was one too long.
mtmd.cpp includes clip-impl.h and not common.h, so that is the overload the `<image: N>`
ordinal label is built with, and mtmd_tokenize_text_internal tokenizes text.data() with
text.length(), all size + 1 bytes of it. Verified with llama-tokenize on the VisionPsy vocab:
`<image: 0>` is [44, 5028, 42, 216, 32, 46], the same bytes plus a trailing NUL are
[44, 5028, 42, 216, 32, 46, 190], so every image in a multi-image prompt carried one garbage
token in a position the checkpoint never saw. common/common.cpp:455 is the same helper written
correctly; this brings clip's copy in line, which also drops the stray byte from every error
message built with it.

The sibling slice-delimiter template trims by hand for this reason
(mtmd.cpp:1234). Registered the coverage that was missing: the base VisionPsy row in tests.sh
now passes two images, since the label only appears above one image and both existing rows
passed a single --image. Confirmed locally, two images still answer "the new york times".

(cherry picked from commit e38010c8ef7af56a0f86cf6ca747969483218c05)

* QVAC-23075 fix: stop letterboxing the VisionPsy refined image

The projector inherited image_pad_rf = PAD_CEIL, so the refined image was aspect-preserved and
centred inside black bars, while the reference stretches straight to the target,
resize(img, [new_h, new_w]) in DynamicResize.forward. PAD_NONE with the existing bicubic algo
is that stretch. It bites whenever the refined aspect ratio differs from the original, which is
the shipped Flash default and needs no flag: 640x480 refines to 1024x512 and PAD_CEIL leaves
170 black columns on each side, a third of the encoded pixels; 1024x768 refines to 1024x1024
and gains 128 black rows top and bottom.

VisionPsy has its own hparams case, so idefics3 is untouched. HF's Idefics3ImageProcessor also
stretches, so it likely wants the same, but that is a pre-existing question and not this PR's.

Measured on Metal against the previous head: the Flash output changes on 960x720, 1024x768 and
640x480, exactly the sizes whose refined aspect ratio differs, and is unchanged on 256x256 and
on every base size, where the refine preserves the aspect ratio and PAD_CEIL was already adding
nothing. Slice and encode counts are identical throughout.

(cherry picked from commit e356a34e53ba125671f11f4d80f02755d72d1c76)

* QVAC-23075 fix: require the slicing sizes whenever VisionPsy loads, flag off included

The positivity guard only ran when no-upscale was on, but the base rule divides by the same two
values. With the flag off, image_size 0 reaches GGML_ASSERT(align_size > 0) in
calc_size_preserved_ratio and aborts the process at the first image instead of failing the
request, and image_longest_edge 0 makes the refined size {0,0}, so the grid is empty and the
model silently receives the overview alone.

Not extended to idefics3, which is where the reported version of this would have gone. The
shipped ggml-org/SmolVLM-500M-Instruct-GGUF mmproj carries no clip.vision.preproc_image_size
at all, so a throw there stops a model that loads today: verified, it failed to load with the
broader check. That model is already overview-only for exactly this reason, 1 image encode for
a 640x488 input, so it now gets a warning that says so rather than a hard failure. Also left
out of the global path, where image_size 0 legally means dynamic sizing per load_hparams' own
sanity check, so a blanket check would reject Qwen-VL.

Verified after the change: SmolVLM loads, warns, still answers "The New York Times"; VisionPsy
base loads and slices the same 17 encodes as before.

(cherry picked from commit 41266e2080a2ad3db898c14c6711b37c8b1a858b)

* QVAC-23075 test: golden tests for the idefics3 refined size and slice grid

The sizing rule decides how many slices an image becomes, and both bugs it carried into review
were invisible to the end-to-end tests, which only read the answer text. It is now a pure
function, mtmd_calc_idefics3_sizing, called by the preprocessor and checked directly by
tests/test-mtmd-preproc-sizing.cpp against values transcribed from the reference
DynamicResize._get_new_hw and evaluated in doubles. The Python used to produce them is in the
test's header comment.

25 cases: the 4:3 class and its transpose, which is what the float32 scale got wrong; small
square, where Flash lands on a single slice and base upscales to the cap; exact power-of-two
ratios, where nothing moved; the repo's own 640x488 test image, whose short side is under one
slice so Flash takes the enlarging branch; extreme aspect ratios; at and above the cap, where
the two rules converge; and a zero-size input. Each case also asserts the invariants the
splitter needs, both sides a whole number of slices and inside [image_size, cap].

Four of the expected slice counts are cross-checked against hardware: they plus one overview
equal the image-encode counts a local Metal run logs for the same inputs.

Confirmed the test catches the regression it was written for. Putting float32 back in
calc_size_preserved_ratio fails 960x720, 720x960 and 1920x1440 with a 4x4 grid where the
reference gives 4x3, while 1440x1080 still passes, which is why the class needs several
members rather than one.

(cherry picked from commit 8ee5417b1cdde90b174e37fbd2c73336ba060e87)

* QVAC-23075 test: pin the projector alias to general.name

The published VisionPsy mmprojs declare clip.projector_type = "custom", which any future model
could also pick, so the alias only resolves when general.name is VisionPsyNano as well. That
second condition is the safety property and it is one && away from being lost, with no visible
symptom until some unrelated "custom" mmproj loads as VisionPsy and preprocesses wrongly.

tests/test-clip-projector-alias.cpp checks the shipped pair resolves, that the same projector
string with any other name, an empty name, or a case-folded name stays unknown, that the right
name behind a different projector string is not the alias either, and that "custom" is absent
from the canonical table so it can only ever be reached through the name-gated path.

(cherry picked from commit ca62e1594f7f3af54cbfbab79bdcc70d253bd4eb)

* QVAC-23075 test: load-time coverage for the preprocessing metadata and the no-upscale override

tests/test-clip-preproc-metadata.cpp generates metadata-only mmprojs with the gguf writer and
runs them through clip_init, so it needs no committed fixture and no model download. That works
because the sizing check sits between load_hparams and load_tensors: a file that passes it still
fails, on a missing tensor, and each positive case asserts that specific failure so a load that
stopped earlier cannot masquerade as a pass.

13 cases. Rejected: image_size 0 and cap 0, each with the flag on and off, since the base rule
divides by the same values; and a cap below one slice, which is the std::clamp precondition.
Accepted: the shipped 512 and 2048 shape in both variants, and a cap of exactly one slice.

The rest covers the override. -1 leaves the GGUF value alone and logs no custom value, 0 and 1
both apply, turning it off against a GGUF that turned it on is announced because that is what a
zero-initialized params struct passes, and a projector that does not read the flag says it is
ignoring it. idefics3 without the cap key keeps loading and warns that slicing is off, which is
the state the shipped ggml-org/SmolVLM-500M-Instruct-GGUF mmproj is in.

(cherry picked from commit d52ab394a9bffd3fe8335d2188142969ec51f40d)

* QVAC-23075 test: check the VisionPsy tile structure in tests.sh, not just the answer

The answer text survives a wrong slice count, which is how the float32 sizing bug and the
letterboxed refine both passed this file. Rows can now declare how many image encodes they
expect, one per slice plus the overview, and the run fails when the count moves.

Both VisionPsy rows declare one. test-1.jpeg is 640x488: the base rule refines it to 2048x2048,
a 4x4 grid plus overview, and the row passes the image twice for the ordinal labels, so 34. The
Flash rule refines the same image to 1024x512, a 2x1 grid plus overview, so 3. Both counts
measured against the published checkpoints on Metal.

Only rows that declare a count are checked, so nothing else in the file changes.

(cherry picked from commit 8be2131e27d5589667db9c7f780fc1a9ef20869a)

* QVAC-23075 fix: reject a zero image_size on idefics3 too, not only VisionPsy

The previous guard was scoped to VisionPsy, so idefics3 kept the half of the finding that
aborts: image_size is the divisor and the align size of the shared rule, and zero reaches
GGML_ASSERT(align_size > 0) in calc_size_preserved_ratio at the first image, which kills the
process instead of failing the request. Nothing about that is VisionPsy specific, so both
projectors now fail the load.

The cap is the only half that legitimately differs. VisionPsy's published mmprojs all carry
clip.vision.preproc_image_size, so a missing cap there is broken metadata and throws, while
ggml-org/SmolVLM-500M-Instruct-GGUF carries none and would stop loading, so idefics3 keeps
the warning that says slicing is off.

Verified that SmolVLM still loads, still warns and still answers "The New York Times" on
tools/mtmd/test-1.jpeg, 1 encode as before. tests/test-clip-preproc-metadata.cpp gains the
idefics3 zero-image-size cases, both flag states, plus a valid idefics3 row so the new throw
cannot start rejecting a good file: 16 cases, all passing.

(cherry picked from commit 55b4a6fdceb73a59d8b94326a50755a3dbae1078)

* QVAC-23075 docs: say that libmtmd consumers must rebuild on every update

The ABI question raised in review: mtmd_context_params is passed and returned by value and
gains fields over time, so its size changes while SOVERSION stays 0. Upstream set SOVERSION 0
in #17091 and has appended fields in #17652, #24384 and #24865 without touching it, so
bumping it here would diverge for a case our consumers do not have, since they all build
fabric from source through a pinned vcpkg port. Documenting the requirement is the part that
costs nothing and is true regardless of how a consumer links.

(cherry picked from commit 24087591c936d03346c104a0fc9b6e82de17d26d)

* QVAC-23075 test: assert the VisionPsy prompt structure, not only the encode count

The QA request was for token and chunk sequence coverage, and the encode count does not give
it: it is blind to the ordinal labels, to the row and column delimiters and to where the
overview sits. tests.sh gains expect_log/expect_no_log over the -v prompt-assembly log, and
both VisionPsy rows now pin the whole sequence: the ordinals on the two-image row, the first
and last delimiter of each grid, the slice grid line, the chunk total, and the overview
position.

The overview delimiters are token ids rather than text, so nothing in the log said where the
overview went. Added the one debug line that says it, which also names the 1x1 case for what
it is, the refined slice standing in for the overview.

Both assertions were mutation tested against the local Metal build. Blanking ord_img_tmpl
leaves 34 encodes and chunk total 69 untouched and is caught only by the ordinal patterns.
Flipping ov_img_first leaves 3 encodes and total 7 untouched, still answers "the new york
times", and is caught only by the overview-position patterns. Chunk total 69 for the base row
is 2 x (ordinal text + overview + 16 x (delimiter + slice)) plus the trailing text chunk, and
7 for the Flash row.

Not covered, and not claimed: a 1x1 grid needs an image whose long side is at most 512 and the
only committed image is 640x488, so the single-tile path is asserted by
test-mtmd-preproc-sizing at the sizing level only.

(cherry picked from commit 295226e2d9067f88acfeb6184561d66e2366fc99)

* QVAC-23075 fix: format string_format into the result, not a vector

The NUL fix broke the GCC build. Returning std::string(buf.data(), size) over a
std::vector<char> keeps `size` live to the end of the function, and GCC then duplicates the
size == 0 path, where the vector is one byte, and reports -Wformat-truncation against it for
any caller whose format string carries a long literal. ubuntu-24.04-arm builds mtmd with
-Werror, so granite-speech.cpp failed on the "feature_layer_" literal.

Formatting straight into the std::string fixes both halves: the terminating NUL lands on the
byte past the end that the string already reserves for it, so it is still not part of the
result, and the destination extent is no longer something GCC can deduce, so the warning has
nothing to fire on. The size == 0 early return removes that path outright.

Token-identical on the case the NUL fix was about: the two-image VisionPsy prompt is 2240
prompt tokens before and after, still 34 image encodes, still answers "the new york times".

(cherry picked from commit d85b1f36711c9561e2a17b95491c1574e6e343d5)

* QVAC-23075 test: skip the two tests that need unexported symbols on Windows DLL builds

test-mtmd-preproc-sizing reaches mtmd_calc_idefics3_sizing and test-clip-preproc-metadata
reaches clip_init, clip_free and clip_log_set_callback. None of them carry MTMD_API, so
lld-link cannot resolve them across the Windows mtmd.dll boundary and every windows job failed
to link. clip.h says this outright at clip_fa_effective_min_kv, and tests/CMakeLists.txt
already guards a block of llama tests the same way, so use that guard rather than exporting
internals or letting the tests decide the library's export surface.

They still run on Linux and macOS, and on Windows in a static build, which is where the guard
condition ends. test-clip-projector-alias stays outside it: alias resolution is inline in
clip-impl.h, so it links everywhere.

(cherry picked from commit 4bccd82aaffcb03ae2ea3d1d0bd52fe10a8e00dd)

* squash! ci : functionally test the OpenCL ops on an Intel CPU ICD

ci : build the opencl test job with GGML_NATIVE=OFF

The ubuntu-24-opencl job flip-flops between pass and SIGILL (exit 132) on
unrelated commits. GGML_NATIVE defaults to ON for a native x86 build, so
-march=native is baked into the CPU backend objects, and the job's ccache
key is shared across runs. An object built on an AVX-512 Azure runner then
gets reused on a runner without it, and test-backend-ops dies with an
illegal instruction before printing any test output.

Every other x64 job in this repo already passes -DGGML_NATIVE=OFF.

(cherry picked from commit aff60f874b1e6520cd7964738f59077a45f13714)

* QVAC-23075 fix: bound the idefics3-style preprocessing metadata at load

The check added for VisionPsy rejected a cap of zero but nothing above it, and
preproc_image_size is a GGUF u32 read into an int. Both sizing rules upscale to
the cap, so the cap alone decides the grid: 512*195 gives a 195x195 grid at one
reserved tile each, and a cap near INT32_MAX overflows the multiply-back in
calc_size_preserved_ratio first. Bound the implied grid against
CLIP_PREPROC_MAX_TILES_LIMIT, the ceiling the Qwen-VL path already clamps to.

A cap that is not a whole number of slices was unchecked too. The slicing loop
steps by image_size and calc_size_no_upscale clamps the long side down to the
cap, so an off-grid cap emits a ragged trailing slice the reference splitter
never produces. test-mtmd-preproc-sizing already asserts that invariant, so
enforce it against real metadata as well.

The idefics3 overview-only warning and the no-upscale rejection were two
independent ifs, so a SmolVLM mmproj with no cap and --image-no-upscale on
logged "slicing is effectively off" and then threw. The later checks are else-if
now, so the warning describes what actually happens.

Only the VisionPsy case read clip.vision.preproc_no_upscale, while the override
and the CLI help both treat idefics3 as accepting the same rule, so an idefics3
GGUF declaring it was silently getting base preprocessing. Read the key there
too.

(cherry picked from commit 4ef2b3fdc0788d38e4e176030c07241ede40c5d0)

* opencl : fix Adreno MoE repack aliasing

Assisted-by: GPT-5.6 Sol
(cherry picked from commit e42aeed9b892e908fa58343a40751141a9d758fb)

* opencl : keep MoE repack words in GEMM order

Assisted-by: GPT-5.6 Sol
(cherry picked from commit 77c2bedd59a1724382d781c02b26d800d95a0a98)

* opencl : add CPU MoE repack diagnostic

Add an environment-gated host repack path to determine whether the Adreno E031.47 compiler still corrupts Q4_K trans4_ns conversion.

Assisted-by: GPT-5.6 Sol
(cherry picked from commit 8c03fb4e41b559f6bc9863fe73cd6b2b6fae099a)

* opencl : default MoE CPU repack on E031.47

Avoid the miscompiled Q4_K trans4 kernel on affected Qualcomm drivers while retaining an environment override for diagnostics.

(cherry picked from commit e36da3e1b63a20153fab7c279d2d3bc054dff909)

* opencl : detect Adreno 830 MoE repack workaround

Use the device name when Android omits the E031.47 compiler token from CL_DRIVER_VERSION so affected phones take the host repack path.

(cherry picked from commit f8749156fe3a6054e6f053191d886d2feb9968b8)

* opencl : disable Q4_K MoE MUL_MAT_ID on Adreno 8xx

The optimized Q4_K MoE kernels produce corrupted results on Adreno 8xx
(Adreno 830, Snapdragon 8 Elite): a MoE model generates garbage tokens
while the same weights are correct on the CPU backend.

Repacking the weights on the host instead of running
kernel_convert_block_q4_k_trans4_ns produces byte-identical corrupt
output, so the weight layout is not at fault - the matmul kernels are.
Declining GGML_OP_MUL_MAT_ID for Q4_K sends the op to the CPU backend,
which is the only configuration that yields correct output on this
hardware. The weight upload keeps using the MoE layout so other Adreno
generations and quantizations are unaffected.

The guard can be lifted per driver with GGML_OPENCL_ADRENO_Q4K_MOE=1.

Also drops the compiler-version heuristic from the host repack helper:
it is a diagnostic aid, not a fix, and stays opt-in via
GGML_OPENCL_Q4K_MOE_CPU_REPACK.

(cherry picked from commit 7f57e21d5edf57a05153cd7717280ef0ccdd0fd1)

* opencl : scope Q4_K MoE guard to the E031.47 compiler

Reading CL_DRIVER_VERSION off an Adreno 830 gives

  OpenCL 3.0 QUALCOMM build: 0800.74 Compiler E031.47.18.51

so the affected shader compiler is identifiable and the fallback does not
have to apply to every Adreno 8xx device indefinitely. Restrict it to
E031 compilers up to 47 and leave newer ones on the optimized kernels,
matching how adreno_e17_compiler_quirks scopes its workaround.

Also records that the dp4a kernels are corrupt on this driver too: they
produce different but equally invalid output, so preferring them is not a
way to keep the work on the GPU.

(cherry picked from commit 1c57effdbdcbc07835834b7f680e5b0977776ab6)

* opencl : extend MoE guard to every affected quantization

The Q4_K-only guard was narrower than the configuration that was
validated. A Samsung Galaxy S25 Ultra run with it still produced garbage:
the model is Q4_K_M, which carries 12 Q6_K and 6 Q5_0 tensors alongside
95 Q4_K ones, so 900 MiB of experts stayed on the optimized Adreno MoE
kernels while only the Q4_K ones moved to the CPU.

All of those quantizations share the same kernels, so decline the whole
Adreno-only MUL_MAT_ID branch instead of a single type. That matches the
configuration validated on an Adreno 830 via GGML_OPENCL_OPFILTER, which
declines every MUL_MAT_ID and produces correct output. Q4_0, Q8_0 and
MXFP4 have general MUL_MAT_ID support, are handled earlier, and keep
running on the GPU.

Renames the escape hatch to GGML_OPENCL_ADRENO_MOE_KERNELS to match the
wider scope.

(cherry picked from commit f761955d4f82d0b8d35e85d3a6043a8c53fd35c6)

* opencl : skip padded MoE output stores

The MoE GEMM kernels pointed padding slots at column 0 of the tile and
relied on writing column 0 last to overwrite them:

    if (idx == 0xFFFFFFFF) idx = src2[block_id_n * TILESIZE_N + 0];
    ...
    barrier(CLK_GLOBAL_MEM_FENCE);
    write_imagef(dst, out_idx[0] + m_offset, reg_c.s0);

dst is an image, and image stores are ordered by CLK_IMAGE_MEM_FENCE, not
CLK_GLOBAL_MEM_FENCE, so a driver is free to reorder or coalesce them. When
that happens column 0 of every padded tile keeps a padding accumulator
(zero) instead of its computed value, which silently drops one token's
expert output per tile. kernel_moe_fill pads every expert's last tile, so
prefill hits this constantly.

Keep the sentinel in out_idx and skip padded slots instead. Every store then
targets a distinct address and needs no ordering at all, and padded tiles
issue fewer image writes. ne01 is a multiple of 32 on this path, so a real
idx * ne01 is even and cannot collide with the odd sentinel.

Also move the tail-row return below the barrier: the global size along dim 0
is rounded up to TILESIZE_M, so returning above it let some work-items skip a
barrier the rest of the workgroup waits on.

Covers the quantizations exercised by Q4_K_M MoE models (Q4_K, Q5_0, Q6_K);
the remaining f32_ns variants share the pattern and follow separately.

(cherry picked from commit b78072f1db1e7065d75893541a4fdb61aec908fd)

* opencl : fix Adreno nibble-packing miscompile in trans4_ns repack

The Adreno E031.47 shader compiler miscompiles the two nibble-packing
helpers shared by every *_trans4_ns weight-repack kernel:

  - narrowing intermediates to uchar locals keeps only the first byte of
    each packed word and drops the rest;
  - the low-nibble path additionally loses the mask on the `<< 4` term, so
    the odd byte's high nibble leaks into the next byte of the word.

Because both helpers are shared, every MoE expert weight repacked for the
optimized kernels was silently corrupted regardless of quantization type,
which is why MoE models emitted garbage on Adreno 8xx. The matmul kernels
were not at fault.

Keep all intermediates in uint registers, mask each byte explicitly in
pack_uchar4, and combine low nibbles as `lo + hi * 16u` so no term can
exceed 8 bits even if a mask is elided.

Verified on an S25 Ultra (Adreno 830) by diffing the convert kernel's four
output buffers against a host reference repack: the quant buffer went from
54.2% of bytes wrong to bit-exact, with d/dm/s unaffected throughout. With
the optimized MoE kernels fully enabled, both a text and an OCR reproducer
now match their CPU baselines.

(cherry picked from commit 8286dfb9c5b447b38d30b663460020be46f513f2)

* opencl : fix get_scale_min_k4 miscompile in MoE matmul kernels

The same E031.47 compiler mishandles get_scale_min_k4() when it returns the
6-bit scale and min through pointers to private scalars, yielding random
noise for Q4_K/Q5_K MUL_MAT_ID. Return both values packed in a single uint
instead.

Takes MUL_MAT_ID failures on Adreno 830 from 76 to 0 in test-backend-ops.

(cherry picked from commit 396759ed796961e97531a62f886e99f2655c4310)

* test : cover Adreno MoE repack shapes in test-backend-ops

Add MUL_MAT_ID repack cases using the expert geometry of a real deployed
model (n_embd=1280, n_ff_exp=896, 6 of 64 experts) across both the GEMV
(decode) and GEMM (prefill) paths, which the existing small shapes missed.

Draw expert ids from a deterministic per-row permutation rather than an
ascending run, so the router table is exercised the way a real router drives
it, and report byte/block detail when a repack readback mismatches.

Stop forcing err() to 1.0 when the readback roundtrip fails: that masked the
real matmul error and hid whether the matmul itself was correct.

(cherry picked from commit 41633de416966cfa1d28015eac3f2252694c2b34)

* opencl : re-enable MoE MUL_MAT_ID on Adreno 8xx

The guard added earlier in this branch declined MUL_MAT_ID for every
quantization reaching the optimized MoE kernels on Adreno 8xx with the
E031.47 compiler, sending those ops to the CPU backend because that was the
only configuration known to produce correct output.

That was a mitigation for the two miscompiles fixed earlier in this branch,
so it is no longer needed. Its rationale was also wrong on the decisive
point: it concluded that host-side repacking reproduced the corruption bit
for bit and therefore the defect lay in the matmul kernels rather than the
weight layout. Host repacking in fact produces correct output, and the
defect was in the trans4_ns repack path.

Keeping the guard would fix the miscompile and then disable the code path it
fixes. Measured on a OnePlus CPH2723 (SM8750, Adreno 830, E031.47) with a
Q4_K_M MoE model and no environment overrides, decode goes from 3.15 to
74.55 tokens per second, a 23.7x speedup, while the generated text stays
byte-identical to the CPU-fallback baseline.

The compiler-version detection the guard used stays: it predates this
branch and backs unrelated feature checks.

(cherry picked from commit 040e0b2c37386936e109a83d526bd95ba52096e0)

* test : fail Adreno repack cases on a broken weight roundtrip

The Adreno repack cases reported a failed store/load roundtrip as an
informational note and returned the matmul error unchanged, so a corrupted
repack could still pass as long as that error stayed under the threshold.

The tolerance was added when the mismatch was believed to be a test-only
artifact of a raw read path. It was not — it was the same packing miscompile
fixed earlier in this branch. With the packing helpers corrected, all 80
adreno_trans4_ns cases roundtrip byte-exact on E031.47, so a mismatch is now
treated as the genuine failure it is.

(cherry picked from commit a8553a69209cba26032a3d39efc503f51a07f28b)

* squash! Vulkan: Add MUL_MAT_MAT and MUL_MAT_VEC support for TQ1

Vulkan: Fix TQ1/2 multi_mat_id pipeline

* ggml-cuda: implement gelu backward

Signed-off-by: makaveli10 <vineet.suryan@collabora.com>

* ggml-cuda: implement l2_norm backward

Signed-off-by: makaveli10 <vineet.suryan@collabora.com>

* ggml-vulkan: fix l2_norm_back for non-contiguous inputs

Signed-off-by: makaveli10 <vineet.suryan@collabora.com>

* ggml : add vector index C API foundation

Assisted-by: GPT-5.5

* ggml : harden vector-index foundation
Fix malformed snapshot handling, reserved padding IDs, finite input validation, score clamping, and PR1 test isolation. Make the vector-index foundation default-off behind a standalone  target.

* ggml : fix vector-index shared build coverage

Export vector-index symbols correctly in shared builds by propagating GGML_SHARED, and enable GGML_VECTOR_INDEX in Linux and Windows shared CI so the library and test target are exercised.

* ggml : harden vector-index snapshot tests
Reject non-zero reserved v1 header bytes, expand malformed snapshot and invalid API coverage, and fix ggml package configuration paths.

* ggml : harden vector index persistence and packaging
Make snapshot writes atomic, tighten API edge-case handling, and add CI coverage for package exports and static consumers.

* ggml : fix vector-index API and package coverage

* ggml : fix vector-index packaging and test coverage

* ggml : renumber vector-index error codes

Assisted-by: GPT-5.5

* ggml: harden vector index foundation
Fix vector-index API semantics, snapshot I/O safety, search ranking, and package smoke checks.

* ggml : fix vector-index snapshot file handling

Avoid GCC attribute warnings from the FILE deleter and keep POSIX snapshot writes restrictive until publish. Preserve existing file modes, apply default creation permissions for new files, and cover the permission behavior in tests.

* ggml-vector-index : amortize add capacity growth
Avoid exact reserve calls during add by growing storage capacity with slack and only reserving the id map when insertion would rehash.

* ggml-vector-index : harden snapshot rename durability

Sync the temporary snapshot again after chmod, fsync the parent directory after rename, and report post-rename sync failures as GGML_VEC_INDEX_E_NOT_DURABLE.

* ggml : reject oversized vector index snapshots on write

* ggml : add vector index snapshot helper header

Assisted-by: GPT-5.5

* ggml : harden vector-index durability and packaging
Fix vector-index snapshot durability on Windows and macOS, tighten byte-span overflow validation, and cover static package consumers in CI.

* ggml : harden vector-index durability and packaging
Fix vector-index snapshot durability on Windows and macOS, tighten byte-span overflow validation, and cover static package consumers in CI.

* tests : cover failed vector-index snapshot overwrite
Add regression coverage that a failed atomic write targeting an existing valid snapshot leaves the snapshot bytes unchanged.

* Fixing metal run issues in CI by now uses the venv Python for both pip installs and the jinja2 import check.

* Revert "Fixing metal run issues in CI by now uses the venv Python for both pip installs and the jinja2 import check."

This reverts commit 7be3d08b1934e0bdb3cc91acb743f252ccf500b8.

Assisted-by: GPT-5.5

* ci : use venv python for model conversion

Run converter installs and imports through the CI virtualenv interpreter so macOS self-hosted jobs do not fall back to a Python without jinja2 or torch.

Assisted-by: GPT-5.5

* ci : revert venv python conversion

Restore the CI script after confirming the Apple GPU failure comes from the self-hosted runner Python version rather than converter interpreter selection.

Assisted-by: GPT-5.5

* ggml : add q4 q8 vector index search

Assisted-by: GPT-5.5

* ggml : restore vector-index utility APIs

Assisted-by: GPT-5.5

* ggml : reject trailing vector-index snapshots

Assisted-…
Sign up for free to join this conversation on GitHub. Already have an account? Sign in to comment

Labels

documentation Improvements or additions to documentation ggml changes relating to the ggml tensor library for machine learning merge ready A maintainer can use this label to indicate that they consider the changes final and ready to merge. SYCL https://en.wikipedia.org/wiki/SYCL - GPU programming language testing Everything test related

Projects

None yet

Development

Successfully merging this pull request may close these issues.

8 participants